poclbm120327.cl 40 KB

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