poclbm120222.cl 40 KB

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