diakgcn120216.cl 31 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653
  1. // DiaKGCN 20-02-2012 - OpenCL kernel by Diapolo
  2. //
  3. // Parts and / or ideas for this kernel are based upon the public-domain poclbm project, the phatk kernel by Phateus and the DiabloMiner kernel by DiabloD3.
  4. // The kernel was rewritten by me (Diapolo) and is still public-domain!
  5. #ifdef VECTORS8
  6. typedef uint8 u;
  7. #elif defined VECTORS4
  8. typedef uint4 u;
  9. #elif defined VECTORS2
  10. typedef uint2 u;
  11. #else
  12. typedef uint u;
  13. #endif
  14. #ifdef BITALIGN
  15. #pragma OPENCL EXTENSION cl_amd_media_ops : enable
  16. #ifdef BFI_INT
  17. #define ch(x, y, z) amd_bytealign(x, y, z)
  18. #define ma(x, y, z) amd_bytealign(z ^ x, y, x)
  19. #else
  20. #define ch(x, y, z) bitselect(z, y, x)
  21. #define ma(z, x, y) bitselect(z, y, z ^ x)
  22. #endif
  23. #else
  24. #define ch(x, y, z) (z ^ (x & (y ^ z)))
  25. #define ma(x, y, z) ((x & z) | (y & (x | z)))
  26. #endif
  27. #define rotr15(n) (rotate(n, 15U) ^ rotate(n, 13U) ^ (n >> 10U))
  28. #define rotr25(n) (rotate(n, 25U) ^ rotate(n, 14U) ^ (n >> 3U))
  29. #define rotr26(n) (rotate(n, 26U) ^ rotate(n, 21U) ^ rotate(n, 7U))
  30. #define rotr30(n) (rotate(n, 30U) ^ rotate(n, 19U) ^ rotate(n, 10U))
  31. __kernel
  32. __attribute__((reqd_work_group_size(WORKSIZE, 1, 1)))
  33. void search(
  34. #ifndef GOFFSET
  35. const u base,
  36. #endif
  37. const uint PreVal0, const uint PreVal4,
  38. const uint H1, const uint D1A, const uint B1, const uint C1,
  39. const uint F1, const uint G1, const uint C1addK5, const uint B1addK6, const uint PreVal0addK7,
  40. const uint W16addK16, const uint W17addK17,
  41. const uint PreW18, const uint PreW19,
  42. const uint W16, const uint W17,
  43. const uint PreW31, const uint PreW32,
  44. const uint state0, const uint state1, const uint state2, const uint state3,
  45. const uint state4, const uint state5, const uint state6, const uint state7,
  46. const uint state0A, const uint state0B,
  47. const uint state1A, const uint state2A, const uint state3A, const uint state4A,
  48. const uint state5A, const uint state6A, const uint state7A,
  49. __global uint * output)
  50. {
  51. u V[8];
  52. u W[16];
  53. #ifdef VECTORS8
  54. #ifdef GOFFSET
  55. const u nonce = ((uint)get_global_id(0) << 3) + (u)(0, 1, 2, 3, 4, 5, 6, 7);
  56. #else
  57. const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0) << 3) + ((uint)get_local_id(0) << 3) + base;
  58. #endif
  59. #elif defined VECTORS4
  60. #ifdef GOFFSET
  61. const u nonce = ((uint)get_global_id(0) << 2) + (u)(0, 1, 2, 3);
  62. #else
  63. const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0) << 2) + ((uint)get_local_id(0) << 2) + base;
  64. #endif
  65. #elif defined VECTORS2
  66. #ifdef GOFFSET
  67. const u nonce = ((uint)get_global_id(0) << 1) + (u)(0, 1);
  68. #else
  69. const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0) << 1) + ((uint)get_local_id(0) << 1) + base;
  70. #endif
  71. #else
  72. #ifdef GOFFSET
  73. const u nonce = (uint)get_global_id(0);
  74. #else
  75. const u nonce = ((uint)get_group_id(0) * (uint)get_local_size(0)) + (uint)get_local_id(0) + base;
  76. #endif
  77. #endif
  78. V[0] = PreVal0 + nonce;
  79. V[1] = B1;
  80. V[2] = C1;
  81. V[3] = D1A;
  82. V[4] = PreVal4 + nonce;
  83. V[5] = F1;
  84. V[6] = G1;
  85. V[7] = H1;
  86. V[7] += V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  87. V[3] = V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  88. V[6] += C1addK5 + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  89. V[2] = C1addK5 + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  90. V[5] += B1addK6 + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  91. V[1] = B1addK6 + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  92. V[4] += PreVal0addK7 + nonce + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  93. V[0] = PreVal0addK7 + nonce + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  94. V[3] += 0xd807aa98 + V[7] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  95. V[7] = 0xd807aa98 + V[7] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  96. V[2] += 0x12835b01 + V[6] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  97. V[6] = 0x12835b01 + V[6] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  98. V[1] += 0x243185be + V[5] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  99. V[5] = 0x243185be + V[5] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  100. V[0] += 0x550c7dc3 + V[4] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  101. V[4] = 0x550c7dc3 + V[4] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  102. V[7] += 0x72be5d74 + V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  103. V[3] = 0x72be5d74 + V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  104. V[6] += 0x80deb1fe + V[2] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  105. V[2] = 0x80deb1fe + V[2] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  106. V[5] += 0x9bdc06a7 + V[1] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  107. V[1] = 0x9bdc06a7 + V[1] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  108. V[4] += 0xc19bf3f4 + V[0] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  109. V[0] = 0xc19bf3f4 + V[0] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  110. V[3] += W16addK16 + V[7] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  111. V[7] = W16addK16 + V[7] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  112. V[2] += W17addK17 + V[6] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  113. V[6] = W17addK17 + V[6] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  114. //----------------------------------------------------------------------------------
  115. #ifdef VECTORS8
  116. W[0] = PreW18 + (u)( rotr25(nonce.s0), rotr25(nonce.s0) ^ 0x2004000, rotr25(nonce.s0) ^ 0x4008000, rotr25(nonce.s0) ^ 0x600C000,
  117. rotr25(nonce.s0) ^ 0x8010000, rotr25(nonce.s0) ^ 0xa014000, rotr25(nonce.s0) ^ 0xc018000, rotr25(nonce.s0) ^ 0xe01c000);
  118. #elif defined VECTORS4
  119. W[0] = PreW18 + (u)(rotr25(nonce.x), rotr25(nonce.x) ^ 0x2004000, rotr25(nonce.x) ^ 0x4008000, rotr25(nonce.x) ^ 0x600C000);
  120. #elif defined VECTORS2
  121. W[0] = PreW18 + (u)(rotr25(nonce.x), rotr25(nonce.x) ^ 0x2004000);
  122. #else
  123. W[0] = PreW18 + rotr25(nonce);
  124. #endif
  125. W[1] = PreW19 + nonce;
  126. W[2] = 0x80000000 + rotr15(W[0]);
  127. W[3] = rotr15(W[1]);
  128. W[4] = 0x00000280 + rotr15(W[2]);
  129. W[5] = W16 + rotr15(W[3]);
  130. W[6] = W17 + rotr15(W[4]);
  131. W[7] = W[0] + rotr15(W[5]);
  132. W[8] = W[1] + rotr15(W[6]);
  133. W[9] = W[2] + rotr15(W[7]);
  134. W[10] = W[3] + rotr15(W[8]);
  135. W[11] = W[4] + rotr15(W[9]);
  136. W[12] = W[5] + 0x00a00055 + rotr15(W[10]);
  137. W[13] = W[6] + PreW31 + rotr15(W[11]);
  138. W[14] = W[7] + PreW32 + rotr15(W[12]);
  139. W[15] = W[8] + W17 + rotr15(W[13]) + rotr25(W[0]);
  140. V[1] += 0x0fc19dc6 + V[5] + W[0] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  141. V[5] = 0x0fc19dc6 + V[5] + W[0] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  142. V[0] += 0x240ca1cc + V[4] + W[1] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  143. V[4] = 0x240ca1cc + V[4] + W[1] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  144. V[7] += 0x2de92c6f + V[3] + W[2] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  145. V[3] = 0x2de92c6f + V[3] + W[2] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  146. V[6] += 0x4a7484aa + V[2] + W[3] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  147. V[2] = 0x4a7484aa + V[2] + W[3] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  148. V[5] += 0x5cb0a9dc + V[1] + W[4] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  149. V[1] = 0x5cb0a9dc + V[1] + W[4] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  150. V[4] += 0x76f988da + V[0] + W[5] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  151. V[0] = 0x76f988da + V[0] + W[5] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  152. V[3] += 0x983e5152 + V[7] + W[6] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  153. V[7] = 0x983e5152 + V[7] + W[6] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  154. V[2] += 0xa831c66d + V[6] + W[7] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  155. V[6] = 0xa831c66d + V[6] + W[7] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  156. V[1] += 0xb00327c8 + V[5] + W[8] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  157. V[5] = 0xb00327c8 + V[5] + W[8] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  158. V[0] += 0xbf597fc7 + V[4] + W[9] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  159. V[4] = 0xbf597fc7 + V[4] + W[9] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  160. V[7] += 0xc6e00bf3 + V[3] + W[10] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  161. V[3] = 0xc6e00bf3 + V[3] + W[10] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  162. V[6] += 0xd5a79147 + V[2] + W[11] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  163. V[2] = 0xd5a79147 + V[2] + W[11] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  164. V[5] += 0x06ca6351 + V[1] + W[12] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  165. V[1] = 0x06ca6351 + V[1] + W[12] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  166. V[4] += 0x14292967 + V[0] + W[13] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  167. V[0] = 0x14292967 + V[0] + W[13] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  168. V[3] += 0x27b70a85 + V[7] + W[14] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  169. V[7] = 0x27b70a85 + V[7] + W[14] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  170. V[2] += 0x2e1b2138 + V[6] + W[15] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  171. V[6] = 0x2e1b2138 + V[6] + W[15] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  172. //----------------------------------------------------------------------------------
  173. W[0] = W[0] + W[9] + rotr15(W[14]) + rotr25( W[1]);
  174. W[1] = W[1] + W[10] + rotr15(W[15]) + rotr25( W[2]);
  175. W[2] = W[2] + W[11] + rotr15( W[0]) + rotr25( W[3]);
  176. W[3] = W[3] + W[12] + rotr15( W[1]) + rotr25( W[4]);
  177. W[4] = W[4] + W[13] + rotr15( W[2]) + rotr25( W[5]);
  178. W[5] = W[5] + W[14] + rotr15( W[3]) + rotr25( W[6]);
  179. W[6] = W[6] + W[15] + rotr15( W[4]) + rotr25( W[7]);
  180. W[7] = W[7] + W[0] + rotr15( W[5]) + rotr25( W[8]);
  181. W[8] = W[8] + W[1] + rotr15( W[6]) + rotr25( W[9]);
  182. W[9] = W[9] + W[2] + rotr15( W[7]) + rotr25(W[10]);
  183. W[10] = W[10] + W[3] + rotr15( W[8]) + rotr25(W[11]);
  184. W[11] = W[11] + W[4] + rotr15( W[9]) + rotr25(W[12]);
  185. W[12] = W[12] + W[5] + rotr15(W[10]) + rotr25(W[13]);
  186. W[13] = W[13] + W[6] + rotr15(W[11]) + rotr25(W[14]);
  187. W[14] = W[14] + W[7] + rotr15(W[12]) + rotr25(W[15]);
  188. W[15] = W[15] + W[8] + rotr15(W[13]) + rotr25( W[0]);
  189. V[1] += 0x4d2c6dfc + V[5] + W[0] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  190. V[5] = 0x4d2c6dfc + V[5] + W[0] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  191. V[0] += 0x53380d13 + V[4] + W[1] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  192. V[4] = 0x53380d13 + V[4] + W[1] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  193. V[7] += 0x650a7354 + V[3] + W[2] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  194. V[3] = 0x650a7354 + V[3] + W[2] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  195. V[6] += 0x766a0abb + V[2] + W[3] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  196. V[2] = 0x766a0abb + V[2] + W[3] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  197. V[5] += 0x81c2c92e + V[1] + W[4] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  198. V[1] = 0x81c2c92e + V[1] + W[4] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  199. V[4] += 0x92722c85 + V[0] + W[5] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  200. V[0] = 0x92722c85 + V[0] + W[5] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  201. V[3] += 0xa2bfe8a1 + V[7] + W[6] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  202. V[7] = 0xa2bfe8a1 + V[7] + W[6] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  203. V[2] += 0xa81a664b + V[6] + W[7] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  204. V[6] = 0xa81a664b + V[6] + W[7] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  205. V[1] += 0xc24b8b70 + V[5] + W[8] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  206. V[5] = 0xc24b8b70 + V[5] + W[8] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  207. V[0] += 0xc76c51a3 + V[4] + W[9] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  208. V[4] = 0xc76c51a3 + V[4] + W[9] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  209. V[7] += 0xd192e819 + V[3] + W[10] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  210. V[3] = 0xd192e819 + V[3] + W[10] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  211. V[6] += 0xd6990624 + V[2] + W[11] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  212. V[2] = 0xd6990624 + V[2] + W[11] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  213. V[5] += 0xf40e3585 + V[1] + W[12] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  214. V[1] = 0xf40e3585 + V[1] + W[12] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  215. V[4] += 0x106aa070 + V[0] + W[13] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  216. V[0] = 0x106aa070 + V[0] + W[13] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  217. V[3] += 0x19a4c116 + V[7] + W[14] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  218. V[7] = 0x19a4c116 + V[7] + W[14] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  219. V[2] += 0x1e376c08 + V[6] + W[15] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  220. V[6] = 0x1e376c08 + V[6] + W[15] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  221. //----------------------------------------------------------------------------------
  222. W[0] = W[0] + W[9] + rotr15(W[14]) + rotr25( W[1]);
  223. W[1] = W[1] + W[10] + rotr15(W[15]) + rotr25( W[2]);
  224. W[2] = W[2] + W[11] + rotr15( W[0]) + rotr25( W[3]);
  225. W[3] = W[3] + W[12] + rotr15( W[1]) + rotr25( W[4]);
  226. W[4] = W[4] + W[13] + rotr15( W[2]) + rotr25( W[5]);
  227. W[5] = W[5] + W[14] + rotr15( W[3]) + rotr25( W[6]);
  228. W[6] = W[6] + W[15] + rotr15( W[4]) + rotr25( W[7]);
  229. W[7] = W[7] + W[0] + rotr15( W[5]) + rotr25( W[8]);
  230. W[8] = W[8] + W[1] + rotr15( W[6]) + rotr25( W[9]);
  231. W[9] = W[9] + W[2] + rotr15( W[7]) + rotr25(W[10]);
  232. W[10] = W[10] + W[3] + rotr15( W[8]) + rotr25(W[11]);
  233. W[11] = W[11] + W[4] + rotr15( W[9]) + rotr25(W[12]);
  234. W[12] = W[12] + W[5] + rotr15(W[10]) + rotr25(W[13]);
  235. W[13] = W[13] + W[6] + rotr15(W[11]) + rotr25(W[14]);
  236. V[1] += 0x2748774c + V[5] + W[0] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  237. V[5] = 0x2748774c + V[5] + W[0] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  238. V[0] += 0x34b0bcb5 + V[4] + W[1] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  239. V[4] = 0x34b0bcb5 + V[4] + W[1] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  240. V[7] += 0x391c0cb3 + V[3] + W[2] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  241. V[3] = 0x391c0cb3 + V[3] + W[2] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  242. V[6] += 0x4ed8aa4a + V[2] + W[3] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  243. V[2] = 0x4ed8aa4a + V[2] + W[3] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  244. V[5] += 0x5b9cca4f + V[1] + W[4] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  245. V[1] = 0x5b9cca4f + V[1] + W[4] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  246. V[4] += 0x682e6ff3 + V[0] + W[5] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  247. V[0] = 0x682e6ff3 + V[0] + W[5] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  248. V[3] += 0x748f82ee + V[7] + W[6] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  249. V[7] = 0x748f82ee + V[7] + W[6] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  250. V[2] += 0x78a5636f + V[6] + W[7] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  251. V[6] = 0x78a5636f + V[6] + W[7] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  252. V[1] += 0x84c87814 + V[5] + W[8] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  253. V[5] = 0x84c87814 + V[5] + W[8] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  254. V[0] += 0x8cc70208 + V[4] + W[9] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  255. V[4] = 0x8cc70208 + V[4] + W[9] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  256. V[7] += 0x90befffa + V[3] + W[10] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  257. V[3] = 0x90befffa + V[3] + W[10] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  258. V[6] += 0xa4506ceb + V[2] + W[11] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  259. V[2] = 0xa4506ceb + V[2] + W[11] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  260. V[5] += 0xbef9a3f7 + V[1] + W[12] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  261. V[1] = 0xbef9a3f7 + V[1] + W[12] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  262. V[4] += 0xc67178f2 + V[0] + W[13] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  263. V[0] = 0xc67178f2 + V[0] + W[13] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  264. //----------------------------------------------------------------------------------
  265. W[0] = state0 + V[0] + rotr25(state1 + V[1]);
  266. W[1] = state1 + V[1] + 0x00a00000 + rotr25(state2 + V[2]);
  267. W[2] = state2 + V[2] + rotr15(W[0]) + rotr25(state3 + V[3]);
  268. W[3] = state3 + V[3] + rotr15(W[1]) + rotr25(state4 + V[4]);
  269. W[4] = state4 + V[4] + rotr15(W[2]) + rotr25(state5 + V[5]);
  270. W[5] = state5 + V[5] + rotr15(W[3]) + rotr25(state6 + V[6]);
  271. W[6] = state6 + V[6] + 0x00000100 + rotr15(W[4]) + rotr25(state7 + V[7]);
  272. W[7] = state7 + V[7] + W[0] + 0x11002000 + rotr15(W[5]);
  273. W[8] = W[1] + 0x80000000 + rotr15(W[6]);
  274. W[9] = W[2] + rotr15(W[7]);
  275. W[10] = W[3] + rotr15(W[8]);
  276. W[11] = W[4] + rotr15(W[9]);
  277. W[12] = W[5] + rotr15(W[10]);
  278. W[13] = W[6] + rotr15(W[11]);
  279. W[14] = W[7] + 0x00400022 + rotr15(W[12]);
  280. W[15] = W[8] + 0x00000100 + rotr15(W[13]) + rotr25(W[0]);
  281. // 0x71374491 + 0x1f83d9ab + state1
  282. const u state1AaddV1 = state1A + V[1];
  283. // 0xb5c0fbcf + 0x9b05688c + state2
  284. const u state2AaddV2 = state2A + V[2];
  285. // 0x510e527f + 0xe9b5dba5 + state3
  286. const u state3AaddV3 = state3A + V[3];
  287. // 0x3956c25b + state4
  288. const u state4AaddV4 = state4A + V[4];
  289. // 0x59f111f1 + state5
  290. const u state5AaddV5 = state5A + V[5];
  291. // 0x923f82a4 + state6
  292. const u state6AaddV6 = state6A + V[6];
  293. // 0xab1c5ed5 + state7
  294. const u state7AaddV7 = state7A + V[7];
  295. // 0x98c7e2a2 + state0
  296. V[3] = state0A + V[0];
  297. // 0xfc08884d + state0
  298. V[7] = state0B + V[0];
  299. V[0] = 0x6a09e667;
  300. V[1] = 0xbb67ae85;
  301. V[2] = 0x3c6ef372;
  302. V[4] = 0x510e527f;
  303. V[5] = 0x9b05688c;
  304. V[6] = 0x1f83d9ab;
  305. V[2] += state1AaddV1 + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  306. V[6] = state1AaddV1 + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  307. V[1] += state2AaddV2 + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  308. V[5] = state2AaddV2 + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  309. V[0] += state3AaddV3 + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  310. V[4] = state3AaddV3 + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  311. V[7] += state4AaddV4 + V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  312. V[3] = state4AaddV4 + V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  313. V[6] += state5AaddV5 + V[2] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  314. V[2] = state5AaddV5 + V[2] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  315. V[5] += state6AaddV6 + V[1] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  316. V[1] = state6AaddV6 + V[1] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  317. V[4] += state7AaddV7 + V[0] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  318. V[0] = state7AaddV7 + V[0] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  319. V[3] += 0x5807aa98 + V[7] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  320. V[7] = 0x5807aa98 + V[7] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  321. V[2] += 0x12835b01 + V[6] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  322. V[6] = 0x12835b01 + V[6] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  323. V[1] += 0x243185be + V[5] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  324. V[5] = 0x243185be + V[5] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  325. V[0] += 0x550c7dc3 + V[4] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  326. V[4] = 0x550c7dc3 + V[4] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  327. V[7] += 0x72be5d74 + V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  328. V[3] = 0x72be5d74 + V[3] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  329. V[6] += 0x80deb1fe + V[2] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  330. V[2] = 0x80deb1fe + V[2] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  331. V[5] += 0x9bdc06a7 + V[1] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  332. V[1] = 0x9bdc06a7 + V[1] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  333. V[4] += 0xc19bf274 + V[0] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  334. V[0] = 0xc19bf274 + V[0] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  335. V[3] += 0xe49b69c1 + V[7] + W[0] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  336. V[7] = 0xe49b69c1 + V[7] + W[0] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  337. V[2] += 0xefbe4786 + V[6] + W[1] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  338. V[6] = 0xefbe4786 + V[6] + W[1] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  339. V[1] += 0x0fc19dc6 + V[5] + W[2] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  340. V[5] = 0x0fc19dc6 + V[5] + W[2] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  341. V[0] += 0x240ca1cc + V[4] + W[3] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  342. V[4] = 0x240ca1cc + V[4] + W[3] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  343. V[7] += 0x2de92c6f + V[3] + W[4] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  344. V[3] = 0x2de92c6f + V[3] + W[4] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  345. V[6] += 0x4a7484aa + V[2] + W[5] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  346. V[2] = 0x4a7484aa + V[2] + W[5] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  347. V[5] += 0x5cb0a9dc + V[1] + W[6] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  348. V[1] = 0x5cb0a9dc + V[1] + W[6] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  349. V[4] += 0x76f988da + V[0] + W[7] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  350. V[0] = 0x76f988da + V[0] + W[7] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  351. V[3] += 0x983e5152 + V[7] + W[8] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  352. V[7] = 0x983e5152 + V[7] + W[8] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  353. V[2] += 0xa831c66d + V[6] + W[9] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  354. V[6] = 0xa831c66d + V[6] + W[9] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  355. V[1] += 0xb00327c8 + V[5] + W[10] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  356. V[5] = 0xb00327c8 + V[5] + W[10] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  357. V[0] += 0xbf597fc7 + V[4] + W[11] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  358. V[4] = 0xbf597fc7 + V[4] + W[11] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  359. V[7] += 0xc6e00bf3 + V[3] + W[12] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  360. V[3] = 0xc6e00bf3 + V[3] + W[12] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  361. V[6] += 0xd5a79147 + V[2] + W[13] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  362. V[2] = 0xd5a79147 + V[2] + W[13] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  363. V[5] += 0x06ca6351 + V[1] + W[14] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  364. V[1] = 0x06ca6351 + V[1] + W[14] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  365. V[4] += 0x14292967 + V[0] + W[15] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  366. V[0] = 0x14292967 + V[0] + W[15] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  367. //----------------------------------------------------------------------------------
  368. W[0] = W[0] + W[9] + rotr15(W[14]) + rotr25( W[1]);
  369. W[1] = W[1] + W[10] + rotr15(W[15]) + rotr25( W[2]);
  370. W[2] = W[2] + W[11] + rotr15( W[0]) + rotr25( W[3]);
  371. W[3] = W[3] + W[12] + rotr15( W[1]) + rotr25( W[4]);
  372. W[4] = W[4] + W[13] + rotr15( W[2]) + rotr25( W[5]);
  373. W[5] = W[5] + W[14] + rotr15( W[3]) + rotr25( W[6]);
  374. W[6] = W[6] + W[15] + rotr15( W[4]) + rotr25( W[7]);
  375. W[7] = W[7] + W[0] + rotr15( W[5]) + rotr25( W[8]);
  376. W[8] = W[8] + W[1] + rotr15( W[6]) + rotr25( W[9]);
  377. W[9] = W[9] + W[2] + rotr15( W[7]) + rotr25(W[10]);
  378. W[10] = W[10] + W[3] + rotr15( W[8]) + rotr25(W[11]);
  379. W[11] = W[11] + W[4] + rotr15( W[9]) + rotr25(W[12]);
  380. W[12] = W[12] + W[5] + rotr15(W[10]) + rotr25(W[13]);
  381. W[13] = W[13] + W[6] + rotr15(W[11]) + rotr25(W[14]);
  382. W[14] = W[14] + W[7] + rotr15(W[12]) + rotr25(W[15]);
  383. W[15] = W[15] + W[8] + rotr15(W[13]) + rotr25( W[0]);
  384. V[3] += 0x27b70a85 + V[7] + W[0] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  385. V[7] = 0x27b70a85 + V[7] + W[0] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  386. V[2] += 0x2e1b2138 + V[6] + W[1] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  387. V[6] = 0x2e1b2138 + V[6] + W[1] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  388. V[1] += 0x4d2c6dfc + V[5] + W[2] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  389. V[5] = 0x4d2c6dfc + V[5] + W[2] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  390. V[0] += 0x53380d13 + V[4] + W[3] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  391. V[4] = 0x53380d13 + V[4] + W[3] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  392. V[7] += 0x650a7354 + V[3] + W[4] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  393. V[3] = 0x650a7354 + V[3] + W[4] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  394. V[6] += 0x766a0abb + V[2] + W[5] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  395. V[2] = 0x766a0abb + V[2] + W[5] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  396. V[5] += 0x81c2c92e + V[1] + W[6] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  397. V[1] = 0x81c2c92e + V[1] + W[6] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  398. V[4] += 0x92722c85 + V[0] + W[7] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  399. V[0] = 0x92722c85 + V[0] + W[7] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  400. V[3] += 0xa2bfe8a1 + V[7] + W[8] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  401. V[7] = 0xa2bfe8a1 + V[7] + W[8] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  402. V[2] += 0xa81a664b + V[6] + W[9] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  403. V[6] = 0xa81a664b + V[6] + W[9] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  404. V[1] += 0xc24b8b70 + V[5] + W[10] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  405. V[5] = 0xc24b8b70 + V[5] + W[10] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  406. V[0] += 0xc76c51a3 + V[4] + W[11] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  407. V[4] = 0xc76c51a3 + V[4] + W[11] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  408. V[7] += 0xd192e819 + V[3] + W[12] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  409. V[3] = 0xd192e819 + V[3] + W[12] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  410. V[6] += 0xd6990624 + V[2] + W[13] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  411. V[2] = 0xd6990624 + V[2] + W[13] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  412. V[5] += 0xf40e3585 + V[1] + W[14] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  413. V[1] = 0xf40e3585 + V[1] + W[14] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  414. V[4] += 0x106aa070 + V[0] + W[15] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  415. V[0] = 0x106aa070 + V[0] + W[15] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  416. //----------------------------------------------------------------------------------
  417. W[0] = W[0] + W[9] + rotr15(W[14]) + rotr25( W[1]);
  418. W[1] = W[1] + W[10] + rotr15(W[15]) + rotr25( W[2]);
  419. W[2] = W[2] + W[11] + rotr15( W[0]) + rotr25( W[3]);
  420. W[3] = W[3] + W[12] + rotr15( W[1]) + rotr25( W[4]);
  421. W[4] = W[4] + W[13] + rotr15( W[2]) + rotr25( W[5]);
  422. W[5] = W[5] + W[14] + rotr15( W[3]) + rotr25( W[6]);
  423. W[6] = W[6] + W[15] + rotr15( W[4]) + rotr25( W[7]);
  424. W[7] = W[7] + W[0] + rotr15( W[5]) + rotr25( W[8]);
  425. W[8] = W[8] + W[1] + rotr15( W[6]) + rotr25( W[9]);
  426. W[9] = W[9] + W[2] + rotr15( W[7]) + rotr25(W[10]);
  427. W[10] = W[10] + W[3] + rotr15( W[8]) + rotr25(W[11]);
  428. W[11] = W[11] + W[4] + rotr15( W[9]) + rotr25(W[12]);
  429. W[12] = W[12] + W[5] + rotr15(W[10]) + rotr25(W[13]);
  430. V[3] += 0x19a4c116 + V[7] + W[0] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  431. V[7] = 0x19a4c116 + V[7] + W[0] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  432. V[2] += 0x1e376c08 + V[6] + W[1] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  433. V[6] = 0x1e376c08 + V[6] + W[1] + ch(V[3], V[4], V[5]) + rotr26(V[3]) + rotr30(V[7]) + ma(V[0], V[1], V[7]);
  434. V[1] += 0x2748774c + V[5] + W[2] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  435. V[5] = 0x2748774c + V[5] + W[2] + ch(V[2], V[3], V[4]) + rotr26(V[2]) + rotr30(V[6]) + ma(V[7], V[0], V[6]);
  436. V[0] += 0x34b0bcb5 + V[4] + W[3] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  437. V[4] = 0x34b0bcb5 + V[4] + W[3] + ch(V[1], V[2], V[3]) + rotr26(V[1]) + rotr30(V[5]) + ma(V[6], V[7], V[5]);
  438. V[7] += 0x391c0cb3 + V[3] + W[4] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  439. V[3] = 0x391c0cb3 + V[3] + W[4] + ch(V[0], V[1], V[2]) + rotr26(V[0]) + rotr30(V[4]) + ma(V[5], V[6], V[4]);
  440. V[6] += 0x4ed8aa4a + V[2] + W[5] + ch(V[7], V[0], V[1]) + rotr26(V[7]);
  441. V[2] = 0x4ed8aa4a + V[2] + W[5] + ch(V[7], V[0], V[1]) + rotr26(V[7]) + rotr30(V[3]) + ma(V[4], V[5], V[3]);
  442. V[5] += 0x5b9cca4f + V[1] + W[6] + ch(V[6], V[7], V[0]) + rotr26(V[6]);
  443. V[1] = 0x5b9cca4f + V[1] + W[6] + ch(V[6], V[7], V[0]) + rotr26(V[6]) + rotr30(V[2]) + ma(V[3], V[4], V[2]);
  444. V[4] += 0x682e6ff3 + V[0] + W[7] + ch(V[5], V[6], V[7]) + rotr26(V[5]);
  445. V[0] = 0x682e6ff3 + V[0] + W[7] + ch(V[5], V[6], V[7]) + rotr26(V[5]) + rotr30(V[1]) + ma(V[2], V[3], V[1]);
  446. V[3] += 0x748f82ee + V[7] + W[8] + ch(V[4], V[5], V[6]) + rotr26(V[4]);
  447. V[7] = 0x748f82ee + V[7] + W[8] + ch(V[4], V[5], V[6]) + rotr26(V[4]) + rotr30(V[0]) + ma(V[1], V[2], V[0]);
  448. V[2] += 0x78a5636f + V[6] + W[9] + ch(V[3], V[4], V[5]) + rotr26(V[3]);
  449. V[1] += 0x84c87814 + V[5] + W[10] + ch(V[2], V[3], V[4]) + rotr26(V[2]);
  450. V[0] += 0x8cc70208 + V[4] + W[11] + ch(V[1], V[2], V[3]) + rotr26(V[1]);
  451. V[7] += V[3] + W[12] + ch(V[0], V[1], V[2]) + rotr26(V[0]);
  452. #define FOUND (0x80)
  453. #define NFLAG (0x7F)
  454. #ifdef VECTORS8
  455. V[7] ^= 0x136032ed;
  456. bool result = any(V[7] == 0);
  457. if (result) {
  458. output[FOUND] = FOUND;
  459. if (!V[7].s0)
  460. output[NFLAG & nonce.s0] = nonce.s0;
  461. if (!V[7].s1)
  462. output[NFLAG & nonce.s1] = nonce.s1;
  463. if (!V[7].s2)
  464. output[NFLAG & nonce.s2] = nonce.s2;
  465. if (!V[7].s3)
  466. output[NFLAG & nonce.s3] = nonce.s3;
  467. if (!V[7].s4)
  468. output[NFLAG & nonce.s4] = nonce.s4;
  469. if (!V[7].s5)
  470. output[NFLAG & nonce.s5] = nonce.s5;
  471. if (!V[7].s6)
  472. output[NFLAG & nonce.s6] = nonce.s6;
  473. if (!V[7].s7)
  474. output[NFLAG & nonce.s7] = nonce.s7;
  475. }
  476. #elif defined VECTORS4
  477. V[7] ^= 0x136032ed;
  478. bool result = any(V[7] == 0);
  479. if (result) {
  480. output[FOUND] = FOUND;
  481. if (!V[7].x)
  482. output[NFLAG & nonce.x] = nonce.x;
  483. if (!V[7].y)
  484. output[NFLAG & nonce.y] = nonce.y;
  485. if (!V[7].z)
  486. output[NFLAG & nonce.z] = nonce.z;
  487. if (!V[7].w)
  488. output[NFLAG & nonce.w] = nonce.w;
  489. }
  490. #elif defined VECTORS2
  491. V[7] ^= 0x136032ed;
  492. bool result = any(V[7] == 0);
  493. if (result) {
  494. output[FOUND] = FOUND;
  495. if (!V[7].x)
  496. output[NFLAG & nonce.x] = nonce.x;
  497. if (!V[7].y)
  498. output[NFLAG & nonce.y] = nonce.y;
  499. }
  500. #else
  501. if (V[7] == 0x136032ed)
  502. output[FOUND] = output[NFLAG & nonce] = nonce;
  503. #endif
  504. }