poclbm130302.cl 43 KB

12345678910111213141516171819202122232425262728293031323334353637383940414243444546474849505152535455565758596061626364656667686970717273747576777879808182838485868788899091929394959697989910010110210310410510610710810911011111211311411511611711811912012112212312412512612712812913013113213313413513613713813914014114214314414514614714814915015115215315415515615715815916016116216316416516616716816917017117217317417517617717817918018118218318418518618718818919019119219319419519619719819920020120220320420520620720820921021121221321421521621721821922022122222322422522622722822923023123223323423523623723823924024124224324424524624724824925025125225325425525625725825926026126226326426526626726826927027127227327427527627727827928028128228328428528628728828929029129229329429529629729829930030130230330430530630730830931031131231331431531631731831932032132232332432532632732832933033133233333433533633733833934034134234334434534634734834935035135235335435535635735835936036136236336436536636736836937037137237337437537637737837938038138238338438538638738838939039139239339439539639739839940040140240340440540640740840941041141241341441541641741841942042142242342442542642742842943043143243343443543643743843944044144244344444544644744844945045145245345445545645745845946046146246346446546646746846947047147247347447547647747847948048148248348448548648748848949049149249349449549649749849950050150250350450550650750850951051151251351451551651751851952052152252352452552652752852953053153253353453553653753853954054154254354454554654754854955055155255355455555655755855956056156256356456556656756856957057157257357457557657757857958058158258358458558658758858959059159259359459559659759859960060160260360460560660760860961061161261361461561661761861962062162262362462562662762862963063163263363463563663763863964064164264364464564664764864965065165265365465565665765865966066166266366466566666766866967067167267367467567667767867968068168268368468568668768868969069169269369469569669769869970070170270370470570670770870971071171271371471571671771871972072172272372472572672772872973073173273373473573673773873974074174274374474574674774874975075175275375475575675775875976076176276376476576676776876977077177277377477577677777877978078178278378478578678778878979079179279379479579679779879980080180280380480580680780880981081181281381481581681781881982082182282382482582682782882983083183283383483583683783883984084184284384484584684784884985085185285385485585685785885986086186286386486586686786886987087187287387487587687787887988088188288388488588688788888989089189289389489589689789889990090190290390490590690790890991091191291391491591691791891992092192292392492592692792892993093193293393493593693793893994094194294394494594694794894995095195295395495595695795895996096196296396496596696796896997097197297397497597697797897998098198298398498598698798898999099199299399499599699799899910001001100210031004100510061007100810091010101110121013101410151016101710181019102010211022102310241025102610271028102910301031103210331034103510361037103810391040104110421043104410451046104710481049105010511052105310541055105610571058105910601061106210631064106510661067106810691070107110721073107410751076107710781079108010811082108310841085108610871088108910901091109210931094109510961097109810991100110111021103110411051106110711081109111011111112111311141115111611171118111911201121112211231124112511261127112811291130113111321133113411351136113711381139114011411142114311441145114611471148114911501151115211531154115511561157115811591160116111621163116411651166116711681169117011711172117311741175117611771178117911801181118211831184118511861187118811891190119111921193119411951196119711981199120012011202120312041205120612071208120912101211121212131214121512161217121812191220122112221223122412251226122712281229123012311232123312341235123612371238123912401241124212431244124512461247124812491250125112521253125412551256125712581259126012611262126312641265126612671268126912701271127212731274127512761277127812791280128112821283128412851286128712881289129012911292129312941295129612971298129913001301130213031304130513061307130813091310131113121313131413151316131713181319132013211322132313241325132613271328132913301331133213331334133513361337133813391340134113421343134413451346134713481349135013511352135313541355135613571358135913601361136213631364136513661367136813691370137113721373137413751376137713781379138013811382138313841385138613871388
  1. // -ck modified kernel taken from Phoenix taken from poclbm, with aspects of
  2. // phatk and others.
  3. // Modified version copyright 2011-2013 Con Kolivas
  4. // This file is taken and modified from the public-domain poclbm project, and
  5. // we have therefore decided to keep it public-domain in Phoenix.
  6. #ifdef VECTORS4
  7. typedef uint4 u;
  8. #elif defined VECTORS2
  9. typedef uint2 u;
  10. #else
  11. typedef uint u;
  12. #endif
  13. __constant uint K[87] = {
  14. 0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
  15. 0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
  16. 0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
  17. 0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
  18. 0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
  19. 0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
  20. 0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
  21. 0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2,
  22. 0xc19bf3f4U,
  23. 0x80000000U,
  24. 0x00000280U,
  25. 0x00a00055U,
  26. 0xf377ed68U,
  27. 0xa54ff53aU,
  28. 0x08909ae5U,
  29. 0x90bb1e3cU,
  30. 0x9b05688cU,
  31. 0xca0b3af3U,
  32. 0x3c6ef372U,
  33. 0xbb67ae85U,
  34. 0x6a09e667U,
  35. 0x50c6645bU,
  36. 0x510e527fU,
  37. 0x3ac42e24U,
  38. 0x5807aa98U,
  39. 0xc19bf274U,
  40. 0x00a00000U,
  41. 0x00000100U,
  42. 0x11002000U,
  43. 0x00400022U,
  44. 0x136032edU
  45. };
  46. #define xc19bf3f4U K[64]
  47. #define x80000000U K[65]
  48. #define x00000280U K[66]
  49. #define x00a00055U K[67]
  50. #define xf377ed68U K[68]
  51. #define xa54ff53aU K[69]
  52. #define x08909ae5U K[70]
  53. #define x90bb1e3cU K[71]
  54. #define x9b05688cU K[72]
  55. #define xca0b3af3U K[73]
  56. #define x3c6ef372U K[74]
  57. #define xbb67ae85U K[75]
  58. #define x6a09e667U K[76]
  59. #define x50c6645bU K[77]
  60. #define x510e527fU K[78]
  61. #define x3ac42e24U K[79]
  62. #define x5807aa98U K[80]
  63. #define xc19bf274U K[81]
  64. #define x00a00000U K[82]
  65. #define x00000100U K[83]
  66. #define x11002000U K[84]
  67. #define x00400022U K[85]
  68. #define x136032edU K[86]
  69. // This part is not from the stock poclbm kernel. It's part of an optimization
  70. // added in the Phoenix Miner.
  71. // Some AMD devices have a BFI_INT opcode, which behaves exactly like the
  72. // SHA-256 ch function, but provides it in exactly one instruction. If
  73. // detected, use it for ch. Otherwise, construct ch out of simpler logical
  74. // primitives.
  75. #ifdef BITALIGN
  76. #pragma OPENCL EXTENSION cl_amd_media_ops : enable
  77. #define rotr(x, y) amd_bitalign((u)x, (u)x, (u)y)
  78. #else
  79. #define rotr(x, y) rotate((u)x, (u)(32 - y))
  80. #endif
  81. #ifdef BFI_INT
  82. // Well, slight problem... It turns out BFI_INT isn't actually exposed to
  83. // OpenCL (or CAL IL for that matter) in any way. However, there is
  84. // a similar instruction, BYTE_ALIGN_INT, which is exposed to OpenCL via
  85. // amd_bytealign, takes the same inputs, and provides the same output.
  86. // We can use that as a placeholder for BFI_INT and have the application
  87. // patch it after compilation.
  88. // This is the BFI_INT function
  89. #define ch(x, y, z) amd_bytealign(x, y, z)
  90. // Ma can also be implemented in terms of BFI_INT...
  91. #define Ma(x, y, z) amd_bytealign( (z^x), (y), (x) )
  92. // AMD's KernelAnalyzer throws errors compiling the kernel if we use
  93. // amd_bytealign on constants with vectors enabled, so we use this to avoid
  94. // problems. (this is used 4 times, and likely optimized out by the compiler.)
  95. #define Ma2(x, y, z) bitselect((u)x, (u)y, (u)z ^ (u)x)
  96. #else // BFI_INT
  97. //GCN actually fails if manually patched with BFI_INT
  98. #define ch(x, y, z) bitselect((u)z, (u)y, (u)x)
  99. #define Ma(x, y, z) bitselect((u)x, (u)y, (u)z ^ (u)x)
  100. #define Ma2(x, y, z) Ma(x, y, z)
  101. #endif
  102. __kernel
  103. __attribute__((vec_type_hint(u)))
  104. __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
  105. void search(const uint state0, const uint state1, const uint state2, const uint state3,
  106. const uint state4, const uint state5, const uint state6, const uint state7,
  107. const uint b1, const uint c1,
  108. const uint f1, const uint g1, const uint h1,
  109. #ifndef GOFFSET
  110. const u base,
  111. #endif
  112. const uint fw0, const uint fw1, const uint fw2, const uint fw3, const uint fw15, const uint fw01r,
  113. const uint D1A, const uint C1addK5, const uint B1addK6,
  114. const uint W16addK16, const uint W17addK17,
  115. const uint PreVal4addT1, const uint Preval0,
  116. volatile __global uint * output)
  117. {
  118. u Vals[24];
  119. u *W = &Vals[8];
  120. #ifdef GOFFSET
  121. const u nonce = (uint)(get_global_id(0));
  122. #else
  123. const u nonce = base + (uint)(get_global_id(0));
  124. #endif
  125. Vals[5]=Preval0;
  126. Vals[5]+=nonce;
  127. Vals[0]=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  128. Vals[0]+=ch(Vals[5],b1,c1);
  129. Vals[0]+=D1A;
  130. Vals[2]=Vals[0];
  131. Vals[2]+=h1;
  132. Vals[1]=PreVal4addT1;
  133. Vals[1]+=nonce;
  134. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  135. Vals[6]=C1addK5;
  136. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  137. Vals[6]+=ch(Vals[2],Vals[5],b1);
  138. Vals[3]=Vals[6];
  139. Vals[3]+=g1;
  140. Vals[0]+=Ma2(g1,Vals[1],f1);
  141. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  142. Vals[6]+=Ma2(f1,Vals[0],Vals[1]);
  143. Vals[7]=B1addK6;
  144. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  145. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  146. Vals[4]=Vals[7];
  147. Vals[4]+=f1;
  148. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  149. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  150. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  151. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  152. Vals[5]+=K[7];
  153. Vals[1]+=Vals[5];
  154. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  155. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  156. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  157. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  158. Vals[2]+=K[8];
  159. Vals[0]+=Vals[2];
  160. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  161. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  162. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  163. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  164. Vals[3]+=K[9];
  165. Vals[6]+=Vals[3];
  166. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  167. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  168. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  169. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  170. Vals[4]+=K[10];
  171. Vals[7]+=Vals[4];
  172. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  173. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  174. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  175. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  176. Vals[1]+=K[11];
  177. Vals[5]+=Vals[1];
  178. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  179. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  180. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  181. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  182. Vals[0]+=K[12];
  183. Vals[2]+=Vals[0];
  184. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  185. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  186. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  187. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  188. Vals[6]+=K[13];
  189. Vals[3]+=Vals[6];
  190. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  191. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  192. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  193. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  194. Vals[7]+=K[14];
  195. Vals[4]+=Vals[7];
  196. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  197. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  198. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  199. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  200. Vals[5]+=xc19bf3f4U;
  201. Vals[1]+=Vals[5];
  202. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  203. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  204. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  205. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  206. Vals[2]+=W16addK16;
  207. Vals[0]+=Vals[2];
  208. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  209. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  210. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  211. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  212. Vals[3]+=W17addK17;
  213. Vals[6]+=Vals[3];
  214. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  215. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  216. W[2]=(rotr(nonce,7)^rotr(nonce,18)^(nonce>>3U));
  217. W[2]+=fw2;
  218. Vals[4]+=W[2];
  219. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  220. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  221. Vals[4]+=K[18];
  222. Vals[7]+=Vals[4];
  223. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  224. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  225. W[3]=nonce;
  226. W[3]+=fw3;
  227. Vals[1]+=W[3];
  228. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  229. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  230. Vals[1]+=K[19];
  231. Vals[5]+=Vals[1];
  232. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  233. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  234. W[4]=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
  235. W[4]+=x80000000U;
  236. Vals[0]+=W[4];
  237. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  238. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  239. Vals[0]+=K[20];
  240. Vals[2]+=Vals[0];
  241. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  242. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  243. W[5]=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
  244. Vals[6]+=W[5];
  245. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  246. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  247. Vals[6]+=K[21];
  248. Vals[3]+=Vals[6];
  249. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  250. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  251. W[6]=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
  252. W[6]+=x00000280U;
  253. Vals[7]+=W[6];
  254. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  255. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  256. Vals[7]+=K[22];
  257. Vals[4]+=Vals[7];
  258. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  259. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  260. W[7]=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
  261. W[7]+=fw0;
  262. Vals[5]+=W[7];
  263. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  264. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  265. Vals[5]+=K[23];
  266. Vals[1]+=Vals[5];
  267. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  268. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  269. W[8]=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
  270. W[8]+=fw1;
  271. Vals[2]+=W[8];
  272. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  273. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  274. Vals[2]+=K[24];
  275. Vals[0]+=Vals[2];
  276. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  277. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  278. W[9]=W[2];
  279. W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
  280. Vals[3]+=W[9];
  281. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  282. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  283. Vals[3]+=K[25];
  284. Vals[6]+=Vals[3];
  285. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  286. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  287. W[10]=W[3];
  288. W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
  289. Vals[4]+=W[10];
  290. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  291. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  292. Vals[4]+=K[26];
  293. Vals[7]+=Vals[4];
  294. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  295. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  296. W[11]=W[4];
  297. W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
  298. Vals[1]+=W[11];
  299. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  300. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  301. Vals[1]+=K[27];
  302. Vals[5]+=Vals[1];
  303. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  304. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  305. W[12]=W[5];
  306. W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
  307. Vals[0]+=W[12];
  308. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  309. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  310. Vals[0]+=K[28];
  311. Vals[2]+=Vals[0];
  312. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  313. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  314. W[13]=W[6];
  315. W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
  316. Vals[6]+=W[13];
  317. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  318. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  319. Vals[6]+=K[29];
  320. Vals[3]+=Vals[6];
  321. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  322. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  323. W[14]=x00a00055U;
  324. W[14]+=W[7];
  325. W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
  326. Vals[7]+=W[14];
  327. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  328. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  329. Vals[7]+=K[30];
  330. Vals[4]+=Vals[7];
  331. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  332. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  333. W[15]=fw15;
  334. W[15]+=W[8];
  335. W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
  336. Vals[5]+=W[15];
  337. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  338. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  339. Vals[5]+=K[31];
  340. Vals[1]+=Vals[5];
  341. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  342. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  343. W[0]=fw01r;
  344. W[0]+=W[9];
  345. W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
  346. Vals[2]+=W[0];
  347. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  348. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  349. Vals[2]+=K[32];
  350. Vals[0]+=Vals[2];
  351. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  352. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  353. W[1]=fw1;
  354. W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
  355. W[1]+=W[10];
  356. W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
  357. Vals[3]+=W[1];
  358. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  359. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  360. Vals[3]+=K[33];
  361. Vals[6]+=Vals[3];
  362. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  363. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  364. W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
  365. W[2]+=W[11];
  366. W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
  367. Vals[4]+=W[2];
  368. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  369. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  370. Vals[4]+=K[34];
  371. Vals[7]+=Vals[4];
  372. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  373. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  374. W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
  375. W[3]+=W[12];
  376. W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
  377. Vals[1]+=W[3];
  378. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  379. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  380. Vals[1]+=K[35];
  381. Vals[5]+=Vals[1];
  382. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  383. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  384. W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
  385. W[4]+=W[13];
  386. W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
  387. Vals[0]+=W[4];
  388. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  389. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  390. Vals[0]+=K[36];
  391. Vals[2]+=Vals[0];
  392. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  393. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  394. W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
  395. W[5]+=W[14];
  396. W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
  397. Vals[6]+=W[5];
  398. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  399. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  400. Vals[6]+=K[37];
  401. Vals[3]+=Vals[6];
  402. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  403. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  404. W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
  405. W[6]+=W[15];
  406. W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
  407. Vals[7]+=W[6];
  408. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  409. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  410. Vals[7]+=K[38];
  411. Vals[4]+=Vals[7];
  412. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  413. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  414. W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
  415. W[7]+=W[0];
  416. W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
  417. Vals[5]+=W[7];
  418. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  419. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  420. Vals[5]+=K[39];
  421. Vals[1]+=Vals[5];
  422. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  423. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  424. W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
  425. W[8]+=W[1];
  426. W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
  427. Vals[2]+=W[8];
  428. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  429. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  430. Vals[2]+=K[40];
  431. Vals[0]+=Vals[2];
  432. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  433. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  434. W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
  435. W[9]+=W[2];
  436. W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
  437. Vals[3]+=W[9];
  438. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  439. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  440. Vals[3]+=K[41];
  441. Vals[6]+=Vals[3];
  442. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  443. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  444. W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
  445. W[10]+=W[3];
  446. W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
  447. Vals[4]+=W[10];
  448. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  449. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  450. Vals[4]+=K[42];
  451. Vals[7]+=Vals[4];
  452. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  453. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  454. W[11]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
  455. W[11]+=W[4];
  456. W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
  457. Vals[1]+=W[11];
  458. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  459. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  460. Vals[1]+=K[43];
  461. Vals[5]+=Vals[1];
  462. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  463. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  464. W[12]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
  465. W[12]+=W[5];
  466. W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
  467. Vals[0]+=W[12];
  468. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  469. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  470. Vals[0]+=K[44];
  471. Vals[2]+=Vals[0];
  472. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  473. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  474. W[13]+=(rotr(W[14],7)^rotr(W[14],18)^(W[14]>>3U));
  475. W[13]+=W[6];
  476. W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
  477. Vals[6]+=W[13];
  478. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  479. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  480. Vals[6]+=K[45];
  481. Vals[3]+=Vals[6];
  482. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  483. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  484. W[14]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
  485. W[14]+=W[7];
  486. W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
  487. Vals[7]+=W[14];
  488. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  489. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  490. Vals[7]+=K[46];
  491. Vals[4]+=Vals[7];
  492. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  493. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  494. W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
  495. W[15]+=W[8];
  496. W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
  497. Vals[5]+=W[15];
  498. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  499. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  500. Vals[5]+=K[47];
  501. Vals[1]+=Vals[5];
  502. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  503. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  504. W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
  505. W[0]+=W[9];
  506. W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
  507. Vals[2]+=W[0];
  508. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  509. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  510. Vals[2]+=K[48];
  511. Vals[0]+=Vals[2];
  512. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  513. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  514. W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
  515. W[1]+=W[10];
  516. W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
  517. Vals[3]+=W[1];
  518. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  519. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  520. Vals[3]+=K[49];
  521. Vals[6]+=Vals[3];
  522. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  523. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  524. W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
  525. W[2]+=W[11];
  526. W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
  527. Vals[4]+=W[2];
  528. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  529. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  530. Vals[4]+=K[50];
  531. Vals[7]+=Vals[4];
  532. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  533. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  534. W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
  535. W[3]+=W[12];
  536. W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
  537. Vals[1]+=W[3];
  538. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  539. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  540. Vals[1]+=K[51];
  541. Vals[5]+=Vals[1];
  542. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  543. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  544. W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
  545. W[4]+=W[13];
  546. W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
  547. Vals[0]+=W[4];
  548. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  549. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  550. Vals[0]+=K[52];
  551. Vals[2]+=Vals[0];
  552. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  553. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  554. W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
  555. W[5]+=W[14];
  556. W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
  557. Vals[6]+=W[5];
  558. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  559. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  560. Vals[6]+=K[53];
  561. Vals[3]+=Vals[6];
  562. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  563. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  564. W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
  565. W[6]+=W[15];
  566. W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
  567. Vals[7]+=W[6];
  568. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  569. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  570. Vals[7]+=K[54];
  571. Vals[4]+=Vals[7];
  572. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  573. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  574. W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
  575. W[7]+=W[0];
  576. W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
  577. Vals[5]+=W[7];
  578. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  579. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  580. Vals[5]+=K[55];
  581. Vals[1]+=Vals[5];
  582. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  583. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  584. W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
  585. W[8]+=W[1];
  586. W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
  587. Vals[2]+=W[8];
  588. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  589. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  590. Vals[2]+=K[56];
  591. Vals[0]+=Vals[2];
  592. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  593. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  594. W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
  595. W[9]+=W[2];
  596. W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
  597. Vals[3]+=W[9];
  598. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  599. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  600. Vals[3]+=K[57];
  601. Vals[6]+=Vals[3];
  602. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  603. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  604. W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
  605. W[10]+=W[3];
  606. W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
  607. Vals[4]+=W[10];
  608. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  609. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  610. Vals[4]+=K[58];
  611. Vals[7]+=Vals[4];
  612. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  613. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  614. W[11]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
  615. W[11]+=W[4];
  616. W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
  617. Vals[1]+=W[11];
  618. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  619. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  620. Vals[1]+=K[59];
  621. Vals[5]+=Vals[1];
  622. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  623. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  624. W[12]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
  625. W[12]+=W[5];
  626. W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
  627. Vals[0]+=W[12];
  628. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  629. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  630. Vals[0]+=K[60];
  631. Vals[2]+=Vals[0];
  632. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  633. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  634. W[13]+=(rotr(W[14],7)^rotr(W[14],18)^(W[14]>>3U));
  635. W[13]+=W[6];
  636. W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
  637. Vals[6]+=W[13];
  638. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  639. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  640. Vals[6]+=K[61];
  641. Vals[3]+=Vals[6];
  642. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  643. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  644. Vals[7]+=W[14];
  645. Vals[7]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
  646. Vals[7]+=W[7];
  647. Vals[7]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
  648. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  649. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  650. Vals[7]+=K[62];
  651. Vals[4]+=Vals[7];
  652. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  653. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  654. Vals[5]+=W[15];
  655. Vals[5]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
  656. Vals[5]+=W[8];
  657. Vals[5]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
  658. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  659. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  660. Vals[5]+=K[63];
  661. Vals[1]+=Vals[5];
  662. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  663. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  664. Vals[5]+=state0;
  665. W[7]=state7;
  666. W[7]+=Vals[2];
  667. Vals[2]=xf377ed68U;
  668. Vals[2]+=Vals[5];
  669. W[0]=Vals[5];
  670. Vals[5]=x6a09e667U;
  671. W[3]=state3;
  672. W[3]+=Vals[0];
  673. Vals[0]=xa54ff53aU;
  674. Vals[0]+=Vals[2];
  675. Vals[2]+=x08909ae5U;
  676. W[6]=state6;
  677. W[6]+=Vals[3];
  678. Vals[3]=x90bb1e3cU;
  679. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  680. Vals[3]+=(x9b05688cU^(Vals[0]&xca0b3af3U));
  681. Vals[7]+=state1;
  682. Vals[3]+=Vals[7];
  683. W[1]=Vals[7];
  684. Vals[7]=xbb67ae85U;
  685. W[2]=state2;
  686. W[2]+=Vals[6];
  687. Vals[6]=x3c6ef372U;
  688. Vals[6]+=Vals[3];
  689. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  690. Vals[3]+=Ma2(Vals[7],Vals[2],Vals[5]);
  691. W[5]=state5;
  692. W[5]+=Vals[4];
  693. Vals[4]=x50c6645bU;
  694. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  695. Vals[4]+=ch(Vals[6],Vals[0],x510e527fU);
  696. Vals[4]+=W[2];
  697. Vals[7]+=Vals[4];
  698. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  699. Vals[4]+=Ma2(Vals[5],Vals[3],Vals[2]);
  700. W[4]=state4;
  701. W[4]+=Vals[1];
  702. Vals[1]=x3ac42e24U;
  703. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  704. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  705. Vals[1]+=W[3];
  706. Vals[5]+=Vals[1];
  707. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  708. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  709. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  710. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  711. Vals[0]+=K[4];
  712. Vals[0]+=W[4];
  713. Vals[2]+=Vals[0];
  714. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  715. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  716. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  717. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  718. Vals[6]+=K[5];
  719. Vals[6]+=W[5];
  720. Vals[3]+=Vals[6];
  721. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  722. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  723. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  724. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  725. Vals[7]+=K[6];
  726. Vals[7]+=W[6];
  727. Vals[4]+=Vals[7];
  728. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  729. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  730. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  731. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  732. Vals[5]+=K[7];
  733. Vals[5]+=W[7];
  734. Vals[1]+=Vals[5];
  735. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  736. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  737. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  738. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  739. Vals[2]+=x5807aa98U;
  740. Vals[0]+=Vals[2];
  741. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  742. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  743. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  744. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  745. Vals[3]+=K[9];
  746. Vals[6]+=Vals[3];
  747. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  748. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  749. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  750. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  751. Vals[4]+=K[10];
  752. Vals[7]+=Vals[4];
  753. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  754. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  755. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  756. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  757. Vals[1]+=K[11];
  758. Vals[5]+=Vals[1];
  759. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  760. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  761. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  762. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  763. Vals[0]+=K[12];
  764. Vals[2]+=Vals[0];
  765. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  766. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  767. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  768. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  769. Vals[6]+=K[13];
  770. Vals[3]+=Vals[6];
  771. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  772. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  773. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  774. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  775. Vals[7]+=K[14];
  776. Vals[4]+=Vals[7];
  777. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  778. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  779. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  780. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  781. Vals[5]+=xc19bf274U;
  782. Vals[1]+=Vals[5];
  783. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  784. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  785. W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
  786. Vals[2]+=W[0];
  787. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  788. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  789. Vals[2]+=K[16];
  790. Vals[0]+=Vals[2];
  791. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  792. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  793. W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
  794. W[1]+=x00a00000U;
  795. Vals[3]+=W[1];
  796. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  797. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  798. Vals[3]+=K[17];
  799. Vals[6]+=Vals[3];
  800. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  801. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  802. W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
  803. W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
  804. Vals[4]+=W[2];
  805. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  806. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  807. Vals[4]+=K[18];
  808. Vals[7]+=Vals[4];
  809. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  810. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  811. W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
  812. W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
  813. Vals[1]+=W[3];
  814. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  815. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  816. Vals[1]+=K[19];
  817. Vals[5]+=Vals[1];
  818. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  819. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  820. W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
  821. W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
  822. Vals[0]+=W[4];
  823. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  824. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  825. Vals[0]+=K[20];
  826. Vals[2]+=Vals[0];
  827. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  828. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  829. W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
  830. W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
  831. Vals[6]+=W[5];
  832. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  833. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  834. Vals[6]+=K[21];
  835. Vals[3]+=Vals[6];
  836. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  837. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  838. W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
  839. W[6]+=x00000100U;
  840. W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
  841. Vals[7]+=W[6];
  842. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  843. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  844. Vals[7]+=K[22];
  845. Vals[4]+=Vals[7];
  846. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  847. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  848. W[7]+=x11002000U;
  849. W[7]+=W[0];
  850. W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
  851. Vals[5]+=W[7];
  852. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  853. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  854. Vals[5]+=K[23];
  855. Vals[1]+=Vals[5];
  856. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  857. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  858. W[8]=x80000000U;
  859. W[8]+=W[1];
  860. W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
  861. Vals[2]+=W[8];
  862. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  863. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  864. Vals[2]+=K[24];
  865. Vals[0]+=Vals[2];
  866. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  867. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  868. W[9]=W[2];
  869. W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
  870. Vals[3]+=W[9];
  871. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  872. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  873. Vals[3]+=K[25];
  874. Vals[6]+=Vals[3];
  875. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  876. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  877. W[10]=W[3];
  878. W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
  879. Vals[4]+=W[10];
  880. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  881. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  882. Vals[4]+=K[26];
  883. Vals[7]+=Vals[4];
  884. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  885. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  886. W[11]=W[4];
  887. W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
  888. Vals[1]+=W[11];
  889. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  890. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  891. Vals[1]+=K[27];
  892. Vals[5]+=Vals[1];
  893. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  894. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  895. W[12]=W[5];
  896. W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
  897. Vals[0]+=W[12];
  898. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  899. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  900. Vals[0]+=K[28];
  901. Vals[2]+=Vals[0];
  902. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  903. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  904. W[13]=W[6];
  905. W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
  906. Vals[6]+=W[13];
  907. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  908. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  909. Vals[6]+=K[29];
  910. Vals[3]+=Vals[6];
  911. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  912. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  913. W[14]=x00400022U;
  914. W[14]+=W[7];
  915. W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
  916. Vals[7]+=W[14];
  917. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  918. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  919. Vals[7]+=K[30];
  920. Vals[4]+=Vals[7];
  921. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  922. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  923. W[15]=x00000100U;
  924. W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
  925. W[15]+=W[8];
  926. W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
  927. Vals[5]+=W[15];
  928. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  929. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  930. Vals[5]+=K[31];
  931. Vals[1]+=Vals[5];
  932. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  933. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  934. W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
  935. W[0]+=W[9];
  936. W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
  937. Vals[2]+=W[0];
  938. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  939. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  940. Vals[2]+=K[32];
  941. Vals[0]+=Vals[2];
  942. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  943. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  944. W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
  945. W[1]+=W[10];
  946. W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
  947. Vals[3]+=W[1];
  948. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  949. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  950. Vals[3]+=K[33];
  951. Vals[6]+=Vals[3];
  952. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  953. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  954. W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
  955. W[2]+=W[11];
  956. W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
  957. Vals[4]+=W[2];
  958. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  959. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  960. Vals[4]+=K[34];
  961. Vals[7]+=Vals[4];
  962. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  963. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  964. W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
  965. W[3]+=W[12];
  966. W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
  967. Vals[1]+=W[3];
  968. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  969. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  970. Vals[1]+=K[35];
  971. Vals[5]+=Vals[1];
  972. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  973. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  974. W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
  975. W[4]+=W[13];
  976. W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
  977. Vals[0]+=W[4];
  978. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  979. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  980. Vals[0]+=K[36];
  981. Vals[2]+=Vals[0];
  982. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  983. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  984. W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
  985. W[5]+=W[14];
  986. W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
  987. Vals[6]+=W[5];
  988. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  989. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  990. Vals[6]+=K[37];
  991. Vals[3]+=Vals[6];
  992. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  993. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  994. W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
  995. W[6]+=W[15];
  996. W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
  997. Vals[7]+=W[6];
  998. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  999. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  1000. Vals[7]+=K[38];
  1001. Vals[4]+=Vals[7];
  1002. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  1003. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  1004. W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
  1005. W[7]+=W[0];
  1006. W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
  1007. Vals[5]+=W[7];
  1008. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  1009. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  1010. Vals[5]+=K[39];
  1011. Vals[1]+=Vals[5];
  1012. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  1013. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  1014. W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
  1015. W[8]+=W[1];
  1016. W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
  1017. Vals[2]+=W[8];
  1018. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  1019. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  1020. Vals[2]+=K[40];
  1021. Vals[0]+=Vals[2];
  1022. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  1023. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  1024. W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
  1025. W[9]+=W[2];
  1026. W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
  1027. Vals[3]+=W[9];
  1028. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  1029. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  1030. Vals[3]+=K[41];
  1031. Vals[6]+=Vals[3];
  1032. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  1033. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  1034. W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
  1035. W[10]+=W[3];
  1036. W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
  1037. Vals[4]+=W[10];
  1038. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  1039. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  1040. Vals[4]+=K[42];
  1041. Vals[7]+=Vals[4];
  1042. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  1043. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  1044. W[11]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
  1045. W[11]+=W[4];
  1046. W[11]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
  1047. Vals[1]+=W[11];
  1048. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  1049. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  1050. Vals[1]+=K[43];
  1051. Vals[5]+=Vals[1];
  1052. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  1053. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  1054. W[12]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
  1055. W[12]+=W[5];
  1056. W[12]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
  1057. Vals[0]+=W[12];
  1058. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  1059. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  1060. Vals[0]+=K[44];
  1061. Vals[2]+=Vals[0];
  1062. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  1063. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  1064. W[13]+=(rotr(W[14],7)^rotr(W[14],18)^(W[14]>>3U));
  1065. W[13]+=W[6];
  1066. W[13]+=(rotr(W[11],17)^rotr(W[11],19)^(W[11]>>10U));
  1067. Vals[6]+=W[13];
  1068. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  1069. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  1070. Vals[6]+=K[45];
  1071. Vals[3]+=Vals[6];
  1072. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  1073. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  1074. W[14]+=(rotr(W[15],7)^rotr(W[15],18)^(W[15]>>3U));
  1075. W[14]+=W[7];
  1076. W[14]+=(rotr(W[12],17)^rotr(W[12],19)^(W[12]>>10U));
  1077. Vals[7]+=W[14];
  1078. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  1079. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  1080. Vals[7]+=K[46];
  1081. Vals[4]+=Vals[7];
  1082. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  1083. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  1084. W[15]+=(rotr(W[0],7)^rotr(W[0],18)^(W[0]>>3U));
  1085. W[15]+=W[8];
  1086. W[15]+=(rotr(W[13],17)^rotr(W[13],19)^(W[13]>>10U));
  1087. Vals[5]+=W[15];
  1088. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  1089. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  1090. Vals[5]+=K[47];
  1091. Vals[1]+=Vals[5];
  1092. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  1093. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  1094. W[0]+=(rotr(W[1],7)^rotr(W[1],18)^(W[1]>>3U));
  1095. W[0]+=W[9];
  1096. W[0]+=(rotr(W[14],17)^rotr(W[14],19)^(W[14]>>10U));
  1097. Vals[2]+=W[0];
  1098. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  1099. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  1100. Vals[2]+=K[48];
  1101. Vals[0]+=Vals[2];
  1102. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  1103. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  1104. W[1]+=(rotr(W[2],7)^rotr(W[2],18)^(W[2]>>3U));
  1105. W[1]+=W[10];
  1106. W[1]+=(rotr(W[15],17)^rotr(W[15],19)^(W[15]>>10U));
  1107. Vals[3]+=W[1];
  1108. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  1109. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  1110. Vals[3]+=K[49];
  1111. Vals[6]+=Vals[3];
  1112. Vals[3]+=(rotr(Vals[2],2)^rotr(Vals[2],13)^rotr(Vals[2],22));
  1113. Vals[3]+=Ma(Vals[7],Vals[2],Vals[5]);
  1114. W[2]+=(rotr(W[3],7)^rotr(W[3],18)^(W[3]>>3U));
  1115. W[2]+=W[11];
  1116. W[2]+=(rotr(W[0],17)^rotr(W[0],19)^(W[0]>>10U));
  1117. Vals[4]+=W[2];
  1118. Vals[4]+=(rotr(Vals[6],6)^rotr(Vals[6],11)^rotr(Vals[6],25));
  1119. Vals[4]+=ch(Vals[6],Vals[0],Vals[1]);
  1120. Vals[4]+=K[50];
  1121. Vals[7]+=Vals[4];
  1122. Vals[4]+=(rotr(Vals[3],2)^rotr(Vals[3],13)^rotr(Vals[3],22));
  1123. Vals[4]+=Ma(Vals[5],Vals[3],Vals[2]);
  1124. W[3]+=(rotr(W[4],7)^rotr(W[4],18)^(W[4]>>3U));
  1125. W[3]+=W[12];
  1126. W[3]+=(rotr(W[1],17)^rotr(W[1],19)^(W[1]>>10U));
  1127. Vals[1]+=W[3];
  1128. Vals[1]+=(rotr(Vals[7],6)^rotr(Vals[7],11)^rotr(Vals[7],25));
  1129. Vals[1]+=ch(Vals[7],Vals[6],Vals[0]);
  1130. Vals[1]+=K[51];
  1131. Vals[5]+=Vals[1];
  1132. Vals[1]+=(rotr(Vals[4],2)^rotr(Vals[4],13)^rotr(Vals[4],22));
  1133. Vals[1]+=Ma(Vals[2],Vals[4],Vals[3]);
  1134. W[4]+=(rotr(W[5],7)^rotr(W[5],18)^(W[5]>>3U));
  1135. W[4]+=W[13];
  1136. W[4]+=(rotr(W[2],17)^rotr(W[2],19)^(W[2]>>10U));
  1137. Vals[0]+=W[4];
  1138. Vals[0]+=(rotr(Vals[5],6)^rotr(Vals[5],11)^rotr(Vals[5],25));
  1139. Vals[0]+=ch(Vals[5],Vals[7],Vals[6]);
  1140. Vals[0]+=K[52];
  1141. Vals[2]+=Vals[0];
  1142. Vals[0]+=(rotr(Vals[1],2)^rotr(Vals[1],13)^rotr(Vals[1],22));
  1143. Vals[0]+=Ma(Vals[3],Vals[1],Vals[4]);
  1144. W[5]+=(rotr(W[6],7)^rotr(W[6],18)^(W[6]>>3U));
  1145. W[5]+=W[14];
  1146. W[5]+=(rotr(W[3],17)^rotr(W[3],19)^(W[3]>>10U));
  1147. Vals[6]+=W[5];
  1148. Vals[6]+=(rotr(Vals[2],6)^rotr(Vals[2],11)^rotr(Vals[2],25));
  1149. Vals[6]+=ch(Vals[2],Vals[5],Vals[7]);
  1150. Vals[6]+=K[53];
  1151. Vals[3]+=Vals[6];
  1152. Vals[6]+=(rotr(Vals[0],2)^rotr(Vals[0],13)^rotr(Vals[0],22));
  1153. Vals[6]+=Ma(Vals[4],Vals[0],Vals[1]);
  1154. W[6]+=(rotr(W[7],7)^rotr(W[7],18)^(W[7]>>3U));
  1155. W[6]+=W[15];
  1156. W[6]+=(rotr(W[4],17)^rotr(W[4],19)^(W[4]>>10U));
  1157. Vals[7]+=W[6];
  1158. Vals[7]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  1159. Vals[7]+=ch(Vals[3],Vals[2],Vals[5]);
  1160. Vals[7]+=K[54];
  1161. Vals[4]+=Vals[7];
  1162. Vals[7]+=(rotr(Vals[6],2)^rotr(Vals[6],13)^rotr(Vals[6],22));
  1163. Vals[7]+=Ma(Vals[1],Vals[6],Vals[0]);
  1164. W[7]+=(rotr(W[8],7)^rotr(W[8],18)^(W[8]>>3U));
  1165. W[7]+=W[0];
  1166. W[7]+=(rotr(W[5],17)^rotr(W[5],19)^(W[5]>>10U));
  1167. Vals[5]+=W[7];
  1168. Vals[5]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  1169. Vals[5]+=ch(Vals[4],Vals[3],Vals[2]);
  1170. Vals[5]+=K[55];
  1171. Vals[1]+=Vals[5];
  1172. Vals[5]+=(rotr(Vals[7],2)^rotr(Vals[7],13)^rotr(Vals[7],22));
  1173. Vals[5]+=Ma(Vals[0],Vals[7],Vals[6]);
  1174. W[8]+=(rotr(W[9],7)^rotr(W[9],18)^(W[9]>>3U));
  1175. W[8]+=W[1];
  1176. W[8]+=(rotr(W[6],17)^rotr(W[6],19)^(W[6]>>10U));
  1177. Vals[2]+=W[8];
  1178. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  1179. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  1180. Vals[2]+=K[56];
  1181. Vals[0]+=Vals[2];
  1182. W[9]+=(rotr(W[10],7)^rotr(W[10],18)^(W[10]>>3U));
  1183. W[9]+=W[2];
  1184. W[9]+=(rotr(W[7],17)^rotr(W[7],19)^(W[7]>>10U));
  1185. Vals[3]+=W[9];
  1186. Vals[3]+=(rotr(Vals[0],6)^rotr(Vals[0],11)^rotr(Vals[0],25));
  1187. Vals[3]+=ch(Vals[0],Vals[1],Vals[4]);
  1188. Vals[3]+=K[57];
  1189. Vals[3]+=Vals[6];
  1190. W[10]+=(rotr(W[11],7)^rotr(W[11],18)^(W[11]>>3U));
  1191. W[10]+=W[3];
  1192. W[10]+=(rotr(W[8],17)^rotr(W[8],19)^(W[8]>>10U));
  1193. Vals[4]+=W[10];
  1194. Vals[4]+=(rotr(Vals[3],6)^rotr(Vals[3],11)^rotr(Vals[3],25));
  1195. Vals[4]+=ch(Vals[3],Vals[0],Vals[1]);
  1196. Vals[4]+=K[58];
  1197. Vals[4]+=Vals[7];
  1198. Vals[1]+=(rotr(Vals[4],6)^rotr(Vals[4],11)^rotr(Vals[4],25));
  1199. Vals[1]+=ch(Vals[4],Vals[3],Vals[0]);
  1200. Vals[1]+=W[11];
  1201. Vals[1]+=(rotr(W[12],7)^rotr(W[12],18)^(W[12]>>3U));
  1202. Vals[1]+=W[4];
  1203. Vals[1]+=(rotr(W[9],17)^rotr(W[9],19)^(W[9]>>10U));
  1204. Vals[1]+=K[59];
  1205. Vals[1]+=Vals[5];
  1206. Vals[2]+=Ma(Vals[6],Vals[5],Vals[7]);
  1207. Vals[2]+=(rotr(Vals[5],2)^rotr(Vals[5],13)^rotr(Vals[5],22));
  1208. Vals[2]+=W[12];
  1209. Vals[2]+=(rotr(W[13],7)^rotr(W[13],18)^(W[13]>>3U));
  1210. Vals[2]+=W[5];
  1211. Vals[2]+=(rotr(W[10],17)^rotr(W[10],19)^(W[10]>>10U));
  1212. Vals[2]+=Vals[0];
  1213. Vals[2]+=(rotr(Vals[1],6)^rotr(Vals[1],11)^rotr(Vals[1],25));
  1214. Vals[2]+=ch(Vals[1],Vals[4],Vals[3]);
  1215. #define FOUND (0x0F)
  1216. #define SETFOUND(Xnonce) output[output[FOUND]++] = Xnonce
  1217. #if defined(VECTORS2) || defined(VECTORS4)
  1218. if (any(Vals[2] == x136032edU)) {
  1219. if (Vals[2].x == x136032edU)
  1220. SETFOUND(nonce.x);
  1221. if (Vals[2].y == x136032edU)
  1222. SETFOUND(nonce.y);
  1223. #if defined(VECTORS4)
  1224. if (Vals[2].z == x136032edU)
  1225. SETFOUND(nonce.z);
  1226. if (Vals[2].w == x136032edU)
  1227. SETFOUND(nonce.w);
  1228. #endif
  1229. }
  1230. #else
  1231. if (Vals[2] == x136032edU)
  1232. SETFOUND(nonce);
  1233. #endif
  1234. }