Guest User

Untitled

a guest
May 15th, 2017
141
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. /**
  2. * Author......: See docs/credits.txt
  3. * License.....: MIT
  4. */
  5.  
  6. #define NEW_SIMD_CODE
  7.  
  8. #include "inc_vendor.cl"
  9. #include "inc_hash_constants.h"
  10. #include "inc_hash_functions.cl"
  11. #include "inc_types.cl"
  12. #include "inc_common.cl"
  13. #include "inc_simd.cl"
  14.  
  15. #if VECT_SIZE == 1
  16. #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)])
  17. #elif VECT_SIZE == 2
  18. #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
  19. #elif VECT_SIZE == 4
  20. #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
  21. #elif VECT_SIZE == 8
  22. #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7])
  23. #elif VECT_SIZE == 16
  24. #define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf])
  25. #endif
  26.  
  27. void m03710m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc)
  28. {
  29. /**
  30. * modifier
  31. */
  32.  
  33. const u32 gid = get_global_id (0);
  34. const u32 lid = get_local_id (0);
  35.  
  36. /**
  37. * salt
  38. */
  39.  
  40. u32 salt_buf0[4];
  41. u32 salt_buf1[4];
  42. u32 salt_buf2[4];
  43. u32 salt_buf3[4];
  44.  
  45. salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
  46. salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
  47. salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
  48. salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
  49. salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
  50. salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
  51. salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
  52. salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
  53. salt_buf2[0] = salt_bufs[salt_pos].salt_buf[ 8];
  54. salt_buf2[1] = salt_bufs[salt_pos].salt_buf[ 9];
  55. salt_buf2[2] = salt_bufs[salt_pos].salt_buf[10];
  56. salt_buf2[3] = salt_bufs[salt_pos].salt_buf[11];
  57. salt_buf3[0] = salt_bufs[salt_pos].salt_buf[12];
  58. salt_buf3[1] = salt_bufs[salt_pos].salt_buf[13];
  59. salt_buf3[2] = salt_bufs[salt_pos].salt_buf[14];
  60. salt_buf3[3] = salt_bufs[salt_pos].salt_buf[15];
  61.  
  62. const u32 salt_len = salt_bufs[salt_pos].salt_len;
  63.  
  64. const u32 pw_salt_len = 16 + salt_len;
  65.  
  66. /**
  67. * loop
  68. */
  69.  
  70. u32 w0l = w0[0];
  71.  
  72. for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
  73. {
  74. const u32x w0r = ix_create_bft (bfs_buf, il_pos);
  75.  
  76. const u32x w0lr = w0l | w0r;
  77.  
  78. u32x w0_t[4];
  79. u32x w1_t[4];
  80. u32x w2_t[4];
  81. u32x w3_t[4];
  82.  
  83. w0_t[0] = w0lr;
  84. w0_t[1] = w0[1];
  85. w0_t[2] = w0[2];
  86. w0_t[3] = w0[3];
  87. w1_t[0] = w1[0];
  88. w1_t[1] = w1[1];
  89. w1_t[2] = w1[2];
  90. w1_t[3] = w1[3];
  91. w2_t[0] = w2[0];
  92. w2_t[1] = w2[1];
  93. w2_t[2] = w2[2];
  94. w2_t[3] = w2[3];
  95. w3_t[0] = w3[0];
  96. w3_t[1] = w3[1];
  97. w3_t[2] = w3[2];
  98. w3_t[3] = w3[3];
  99.  
  100. /**
  101. * md5
  102. */
  103.  
  104. u32x a = MD5M_A;
  105. u32x b = MD5M_B;
  106. u32x c = MD5M_C;
  107. u32x d = MD5M_D;
  108.  
  109. MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
  110. MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
  111. MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
  112. MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
  113. MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
  114. MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
  115. MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
  116. MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
  117. MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
  118. MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
  119. MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
  120. MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
  121. MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
  122. MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
  123. MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
  124. MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
  125.  
  126. MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
  127. MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
  128. MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
  129. MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
  130. MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
  131. MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
  132. MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
  133. MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
  134. MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
  135. MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
  136. MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
  137. MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
  138. MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
  139. MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
  140. MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
  141. MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
  142.  
  143. MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
  144. MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
  145. MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
  146. MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
  147. MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
  148. MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
  149. MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
  150. MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
  151. MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
  152. MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
  153. MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
  154. MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
  155. MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
  156. MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
  157. MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
  158. MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
  159.  
  160. MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
  161. MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
  162. MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
  163. MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
  164. MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
  165. MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
  166. MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
  167. MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
  168. MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
  169. MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
  170. MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
  171. MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
  172. MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
  173. MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
  174. MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
  175. MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
  176.  
  177. a += MD5M_A;
  178. b += MD5M_B;
  179. c += MD5M_C;
  180. d += MD5M_D;
  181.  
  182. w0_t[0] = ((a >> 0) & 255) << 0
  183. | ((a >> 8) & 255) << 16;
  184. w0_t[1] = ((a >> 16) & 255) << 0
  185. | ((a >> 24) & 255) << 16;
  186. w0_t[2] = ((b >> 0) & 255) << 0
  187. | ((b >> 8) & 255) << 16;
  188. w0_t[3] = ((b >> 16) & 255) << 0
  189. | ((b >> 24) & 255) << 16;
  190. w1_t[0] = ((c >> 0) & 255) << 0
  191. | ((c >> 8) & 255) << 16;
  192. w1_t[1] = ((c >> 16) & 255) << 0
  193. | ((c >> 24) & 255) << 16;
  194. w1_t[2] = ((d >> 0) & 255) << 0
  195. | ((d >> 8) & 255) << 16;
  196. w1_t[3] = ((d >> 16) & 255) << 0
  197. | ((d >> 24) & 255) << 16;
  198. w2_t[0] = 0x80;
  199. w2_t[1] = 0;
  200. w2_t[2] = 0;
  201. w2_t[3] = 0;
  202. w3_t[0] = 0;
  203. w3_t[1] = 0;
  204. w3_t[2] = 0;
  205. w3_t[3] = 0;
  206.  
  207. /**
  208. * prepend salt
  209. */
  210.  
  211. switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, salt_len);
  212.  
  213. w3_t[2] = pw_salt_len * 8;
  214. w3_t[3] = 0;
  215.  
  216. w0_t[0] |= a;
  217. w0_t[1] |= b;
  218. w0_t[2] |= c;
  219. w0_t[3] |= d;
  220. w1_t[0] |= 0x80;
  221. w1_t[1] |= 0;
  222. w1_t[2] |= 0;
  223. w1_t[3] |= 0;
  224. w2_t[0] |= 0;
  225. w2_t[1] |= 0;
  226. w2_t[2] |= 0;
  227. w2_t[3] |= 0;
  228. w3_t[0] |= 0;
  229. w3_t[1] |= 0;
  230. w3_t[2] |= 0;
  231. w3_t[3] |= 0;
  232.  
  233. /**
  234. * md5
  235. */
  236.  
  237. a = MD5M_A;
  238. b = MD5M_B;
  239. c = MD5M_C;
  240. d = MD5M_D;
  241.  
  242. MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
  243. MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
  244. MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
  245. MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
  246. MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
  247. MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
  248. MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
  249. MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
  250. MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
  251. MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
  252. MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
  253. MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
  254. MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
  255. MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
  256. MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
  257. MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
  258.  
  259. MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
  260. MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
  261. MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
  262. MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
  263. MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
  264. MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
  265. MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
  266. MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
  267. MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
  268. MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
  269. MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
  270. MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
  271. MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
  272. MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
  273. MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
  274. MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
  275.  
  276. MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
  277. MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
  278. MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
  279. MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
  280. MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
  281. MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
  282. MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
  283. MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
  284. MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
  285. MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
  286. MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
  287. MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
  288. MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
  289. MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
  290. MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
  291. MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
  292.  
  293. MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
  294. MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
  295. MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
  296. MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
  297. MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
  298. MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
  299. MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
  300. MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
  301. MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
  302. MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
  303. MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
  304. MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
  305. MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
  306. MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
  307. MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
  308. MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
  309.  
  310. COMPARE_M_SIMD (a, d, c, b);
  311. }
  312. }
  313.  
  314. void m03710s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, __local u32 *l_bin2asc)
  315. {
  316. /**
  317. * modifier
  318. */
  319.  
  320. const u32 gid = get_global_id (0);
  321. const u32 lid = get_local_id (0);
  322.  
  323. /**
  324. * salt
  325. */
  326.  
  327. u32 salt_buf0[4];
  328. u32 salt_buf1[4];
  329. u32 salt_buf2[4];
  330. u32 salt_buf3[4];
  331.  
  332. salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
  333. salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
  334. salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
  335. salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
  336. salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
  337. salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
  338. salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
  339. salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
  340. salt_buf2[0] = salt_bufs[salt_pos].salt_buf[ 8];
  341. salt_buf2[1] = salt_bufs[salt_pos].salt_buf[ 9];
  342. salt_buf2[2] = salt_bufs[salt_pos].salt_buf[10];
  343. salt_buf2[3] = salt_bufs[salt_pos].salt_buf[11];
  344. salt_buf3[0] = salt_bufs[salt_pos].salt_buf[12];
  345. salt_buf3[1] = salt_bufs[salt_pos].salt_buf[13];
  346. salt_buf3[2] = salt_bufs[salt_pos].salt_buf[14];
  347. salt_buf3[3] = salt_bufs[salt_pos].salt_buf[15];
  348.  
  349. const u32 salt_len = salt_bufs[salt_pos].salt_len;
  350.  
  351. const u32 pw_salt_len = 16 + salt_len;
  352.  
  353. /**
  354. * digest
  355. */
  356.  
  357. const u32 search[4] =
  358. {
  359. digests_buf[digests_offset].digest_buf[DGST_R0],
  360. digests_buf[digests_offset].digest_buf[DGST_R1],
  361. digests_buf[digests_offset].digest_buf[DGST_R2],
  362. digests_buf[digests_offset].digest_buf[DGST_R3]
  363. };
  364.  
  365. /**
  366. * loop
  367. */
  368.  
  369. u32 w0l = w0[0];
  370.  
  371. for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
  372. {
  373. const u32x w0r = ix_create_bft (bfs_buf, il_pos);
  374.  
  375. const u32x w0lr = w0l | w0r;
  376.  
  377. u32x w0_t[4];
  378. u32x w1_t[4];
  379. u32x w2_t[4];
  380. u32x w3_t[4];
  381.  
  382. w0_t[0] = w0lr;
  383. w0_t[1] = w0[1];
  384. w0_t[2] = w0[2];
  385. w0_t[3] = w0[3];
  386. w1_t[0] = w1[0];
  387. w1_t[1] = w1[1];
  388. w1_t[2] = w1[2];
  389. w1_t[3] = w1[3];
  390. w2_t[0] = w2[0];
  391. w2_t[1] = w2[1];
  392. w2_t[2] = w2[2];
  393. w2_t[3] = w2[3];
  394. w3_t[0] = w3[0];
  395. w3_t[1] = w3[1];
  396. w3_t[2] = w3[2];
  397. w3_t[3] = w3[3];
  398.  
  399. /**
  400. * md5
  401. */
  402.  
  403. u32x a = MD5M_A;
  404. u32x b = MD5M_B;
  405. u32x c = MD5M_C;
  406. u32x d = MD5M_D;
  407.  
  408. MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
  409. MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
  410. MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
  411. MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
  412. MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
  413. MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
  414. MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
  415. MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
  416. MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
  417. MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
  418. MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
  419. MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
  420. MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
  421. MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
  422. MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
  423. MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
  424.  
  425. MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
  426. MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
  427. MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
  428. MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
  429. MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
  430. MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
  431. MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
  432. MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
  433. MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
  434. MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
  435. MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
  436. MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
  437. MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
  438. MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
  439. MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
  440. MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
  441.  
  442. MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
  443. MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
  444. MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
  445. MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
  446. MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
  447. MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
  448. MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
  449. MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
  450. MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
  451. MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
  452. MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
  453. MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
  454. MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
  455. MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
  456. MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
  457. MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
  458.  
  459. MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
  460. MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
  461. MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
  462. MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
  463. MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
  464. MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
  465. MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
  466. MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
  467. MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
  468. MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
  469. MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
  470. MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
  471. MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
  472. MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
  473. MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
  474. MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
  475.  
  476. a += MD5M_A;
  477. b += MD5M_B;
  478. c += MD5M_C;
  479. d += MD5M_D;
  480.  
  481. w0_t[0] = ((a >> 0) & 255) << 0
  482. | ((a >> 8) & 255) << 16;
  483. w0_t[1] = ((a >> 16) & 255) << 0
  484. | ((a >> 24) & 255) << 16;
  485. w0_t[2] = ((b >> 0) & 255) << 0
  486. | ((b >> 8) & 255) << 16;
  487. w0_t[3] = ((b >> 16) & 255) << 0
  488. | ((b >> 24) & 255) << 16;
  489. w1_t[0] = ((c >> 0) & 255) << 0
  490. | ((c >> 8) & 255) << 16;
  491. w1_t[1] = ((c >> 16) & 255) << 0
  492. | ((c >> 24) & 255) << 16;
  493. w1_t[2] = ((d >> 0) & 255) << 0
  494. | ((d >> 8) & 255) << 16;
  495. w1_t[3] = ((d >> 16) & 255) << 0
  496. | ((d >> 24) & 255) << 16;
  497. w2_t[0] = 0x80;
  498. w2_t[1] = 0;
  499. w2_t[2] = 0;
  500. w2_t[3] = 0;
  501. w3_t[0] = 0;
  502. w3_t[1] = 0;
  503. w3_t[2] = 0;
  504. w3_t[3] = 0;
  505.  
  506. /**
  507. * prepend salt
  508. */
  509.  
  510. switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, salt_len);
  511.  
  512. w3_t[2] = pw_salt_len * 8;
  513. w3_t[3] = 0;
  514.  
  515. w0_t[0] |= a;
  516. w0_t[1] |= b;
  517. w0_t[2] |= c;
  518. w0_t[3] |= d;
  519. w1_t[0] |= 0x80;
  520. w1_t[1] |= 0;
  521. w1_t[2] |= 0;
  522. w1_t[3] |= 0;
  523. w2_t[0] |= 0;
  524. w2_t[1] |= 0;
  525. w2_t[2] |= 0;
  526. w2_t[3] |= 0;
  527. w3_t[0] |= 0;
  528. w3_t[1] |= 0;
  529. w3_t[2] |= 0;
  530. w3_t[3] |= 0;
  531.  
  532. /**
  533. * md5
  534. */
  535.  
  536. a = MD5M_A;
  537. b = MD5M_B;
  538. c = MD5M_C;
  539. d = MD5M_D;
  540.  
  541. MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
  542. MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
  543. MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
  544. MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
  545. MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
  546. MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
  547. MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
  548. MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
  549. MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
  550. MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
  551. MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
  552. MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
  553. MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
  554. MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
  555. MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
  556. MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
  557.  
  558. MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
  559. MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
  560. MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
  561. MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
  562. MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
  563. MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
  564. MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
  565. MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
  566. MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
  567. MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
  568. MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
  569. MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
  570. MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
  571. MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
  572. MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
  573. MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
  574.  
  575. MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
  576. MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
  577. MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
  578. MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
  579. MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
  580. MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
  581. MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
  582. MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
  583. MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
  584. MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
  585. MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
  586. MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
  587. MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
  588. MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
  589. MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
  590. MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
  591.  
  592. MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
  593. MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
  594. MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
  595. MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
  596. MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
  597. MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
  598. MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
  599. MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
  600. MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
  601. MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
  602. MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
  603. MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
  604. MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
  605.  
  606. if (MATCHES_NONE_VS (a, search[0])) continue;
  607.  
  608. MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
  609. MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
  610. MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
  611.  
  612. COMPARE_S_SIMD (a, d, c, b);
  613. }
  614. }
  615.  
  616. __kernel void m03710_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  617. {
  618. /**
  619. * base
  620. */
  621.  
  622. const u32 gid = get_global_id (0);
  623. const u32 lid = get_local_id (0);
  624. const u32 lsz = get_local_size (0);
  625.  
  626. /**
  627. * bin2asc table
  628. */
  629.  
  630. __local u32 l_bin2asc[256];
  631.  
  632. for (u32 i = lid; i < 256; i += lsz)
  633. {
  634. const u32 i0 = (i >> 0) & 15;
  635. const u32 i1 = (i >> 4) & 15;
  636.  
  637. l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
  638. | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
  639. }
  640.  
  641. barrier (CLK_LOCAL_MEM_FENCE);
  642.  
  643. if (gid >= gid_max) return;
  644.  
  645. /**
  646. * modifier
  647. */
  648.  
  649. u32 w0[4];
  650.  
  651. w0[0] = pws[gid].i[ 0];
  652. w0[1] = pws[gid].i[ 1];
  653. w0[2] = pws[gid].i[ 2];
  654. w0[3] = pws[gid].i[ 3];
  655.  
  656. u32 w1[4];
  657.  
  658. w1[0] = 0;
  659. w1[1] = 0;
  660. w1[2] = 0;
  661. w1[3] = 0;
  662.  
  663. u32 w2[4];
  664.  
  665. w2[0] = 0;
  666. w2[1] = 0;
  667. w2[2] = 0;
  668. w2[3] = 0;
  669.  
  670. u32 w3[4];
  671.  
  672. w3[0] = 0;
  673. w3[1] = 0;
  674. w3[2] = pws[gid].i[14];
  675. w3[3] = 0;
  676.  
  677. const u32 pw_len = pws[gid].pw_len;
  678.  
  679. /**
  680. * main
  681. */
  682.  
  683. m03710m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, l_bin2asc);
  684. }
  685.  
  686. __kernel void m03710_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  687. {
  688. /**
  689. * base
  690. */
  691.  
  692. const u32 gid = get_global_id (0);
  693. const u32 lid = get_local_id (0);
  694. const u32 lsz = get_local_size (0);
  695.  
  696. /**
  697. * modifier
  698. */
  699.  
  700. u32 w0[4];
  701.  
  702. w0[0] = pws[gid].i[ 0];
  703. w0[1] = pws[gid].i[ 1];
  704. w0[2] = pws[gid].i[ 2];
  705. w0[3] = pws[gid].i[ 3];
  706.  
  707. u32 w1[4];
  708.  
  709. w1[0] = pws[gid].i[ 4];
  710. w1[1] = pws[gid].i[ 5];
  711. w1[2] = pws[gid].i[ 6];
  712. w1[3] = pws[gid].i[ 7];
  713.  
  714. u32 w2[4];
  715.  
  716. w2[0] = 0;
  717. w2[1] = 0;
  718. w2[2] = 0;
  719. w2[3] = 0;
  720.  
  721. u32 w3[4];
  722.  
  723. w3[0] = 0;
  724. w3[1] = 0;
  725. w3[2] = pws[gid].i[14];
  726. w3[3] = 0;
  727.  
  728. const u32 pw_len = pws[gid].pw_len;
  729.  
  730. /**
  731. * bin2asc table
  732. */
  733.  
  734. __local u32 l_bin2asc[256];
  735.  
  736. for (u32 i = lid; i < 256; i += lsz)
  737. {
  738. const u32 i0 = (i >> 0) & 15;
  739. const u32 i1 = (i >> 4) & 15;
  740.  
  741. l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
  742. | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
  743. }
  744.  
  745. barrier (CLK_LOCAL_MEM_FENCE);
  746.  
  747. if (gid >= gid_max) return;
  748.  
  749. /**
  750. * main
  751. */
  752.  
  753. m03710m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, l_bin2asc);
  754. }
  755.  
  756. __kernel void m03710_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  757. {
  758. /**
  759. * base
  760. */
  761.  
  762. const u32 gid = get_global_id (0);
  763. const u32 lid = get_local_id (0);
  764. const u32 lsz = get_local_size (0);
  765.  
  766. /**
  767. * modifier
  768. */
  769.  
  770. u32 w0[4];
  771.  
  772. w0[0] = pws[gid].i[ 0];
  773. w0[1] = pws[gid].i[ 1];
  774. w0[2] = pws[gid].i[ 2];
  775. w0[3] = pws[gid].i[ 3];
  776.  
  777. u32 w1[4];
  778.  
  779. w1[0] = pws[gid].i[ 4];
  780. w1[1] = pws[gid].i[ 5];
  781. w1[2] = pws[gid].i[ 6];
  782. w1[3] = pws[gid].i[ 7];
  783.  
  784. u32 w2[4];
  785.  
  786. w2[0] = pws[gid].i[ 8];
  787. w2[1] = pws[gid].i[ 9];
  788. w2[2] = pws[gid].i[10];
  789. w2[3] = pws[gid].i[11];
  790.  
  791. u32 w3[4];
  792.  
  793. w3[0] = pws[gid].i[12];
  794. w3[1] = pws[gid].i[13];
  795. w3[2] = pws[gid].i[14];
  796. w3[3] = pws[gid].i[15];
  797.  
  798. const u32 pw_len = pws[gid].pw_len;
  799.  
  800. /**
  801. * bin2asc table
  802. */
  803.  
  804. __local u32 l_bin2asc[256];
  805.  
  806. for (u32 i = lid; i < 256; i += lsz)
  807. {
  808. const u32 i0 = (i >> 0) & 15;
  809. const u32 i1 = (i >> 4) & 15;
  810.  
  811. l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
  812. | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
  813. }
  814.  
  815. barrier (CLK_LOCAL_MEM_FENCE);
  816.  
  817. if (gid >= gid_max) return;
  818.  
  819. /**
  820. * main
  821. */
  822.  
  823. m03710m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, l_bin2asc);
  824. }
  825.  
  826. __kernel void m03710_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  827. {
  828. /**
  829. * base
  830. */
  831.  
  832. const u32 gid = get_global_id (0);
  833. const u32 lid = get_local_id (0);
  834. const u32 lsz = get_local_size (0);
  835.  
  836. /**
  837. * modifier
  838. */
  839.  
  840. u32 w0[4];
  841.  
  842. w0[0] = pws[gid].i[ 0];
  843. w0[1] = pws[gid].i[ 1];
  844. w0[2] = pws[gid].i[ 2];
  845. w0[3] = pws[gid].i[ 3];
  846.  
  847. u32 w1[4];
  848.  
  849. w1[0] = 0;
  850. w1[1] = 0;
  851. w1[2] = 0;
  852. w1[3] = 0;
  853.  
  854. u32 w2[4];
  855.  
  856. w2[0] = 0;
  857. w2[1] = 0;
  858. w2[2] = 0;
  859. w2[3] = 0;
  860.  
  861. u32 w3[4];
  862.  
  863. w3[0] = 0;
  864. w3[1] = 0;
  865. w3[2] = pws[gid].i[14];
  866. w3[3] = 0;
  867.  
  868. const u32 pw_len = pws[gid].pw_len;
  869.  
  870. /**
  871. * bin2asc table
  872. */
  873.  
  874. __local u32 l_bin2asc[256];
  875.  
  876. for (u32 i = lid; i < 256; i += lsz)
  877. {
  878. const u32 i0 = (i >> 0) & 15;
  879. const u32 i1 = (i >> 4) & 15;
  880.  
  881. l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
  882. | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
  883. }
  884.  
  885. barrier (CLK_LOCAL_MEM_FENCE);
  886.  
  887. if (gid >= gid_max) return;
  888.  
  889. /**
  890. * main
  891. */
  892.  
  893. m03710s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, l_bin2asc);
  894. }
  895.  
  896. __kernel void m03710_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  897. {
  898. /**
  899. * base
  900. */
  901.  
  902. const u32 gid = get_global_id (0);
  903. const u32 lid = get_local_id (0);
  904. const u32 lsz = get_local_size (0);
  905.  
  906. /**
  907. * modifier
  908. */
  909.  
  910. u32 w0[4];
  911.  
  912. w0[0] = pws[gid].i[ 0];
  913. w0[1] = pws[gid].i[ 1];
  914. w0[2] = pws[gid].i[ 2];
  915. w0[3] = pws[gid].i[ 3];
  916.  
  917. u32 w1[4];
  918.  
  919. w1[0] = pws[gid].i[ 4];
  920. w1[1] = pws[gid].i[ 5];
  921. w1[2] = pws[gid].i[ 6];
  922. w1[3] = pws[gid].i[ 7];
  923.  
  924. u32 w2[4];
  925.  
  926. w2[0] = 0;
  927. w2[1] = 0;
  928. w2[2] = 0;
  929. w2[3] = 0;
  930.  
  931. u32 w3[4];
  932.  
  933. w3[0] = 0;
  934. w3[1] = 0;
  935. w3[2] = pws[gid].i[14];
  936. w3[3] = 0;
  937.  
  938. const u32 pw_len = pws[gid].pw_len;
  939.  
  940. /**
  941. * bin2asc table
  942. */
  943.  
  944. __local u32 l_bin2asc[256];
  945.  
  946. for (u32 i = lid; i < 256; i += lsz)
  947. {
  948. const u32 i0 = (i >> 0) & 15;
  949. const u32 i1 = (i >> 4) & 15;
  950.  
  951. l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
  952. | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
  953. }
  954.  
  955. barrier (CLK_LOCAL_MEM_FENCE);
  956.  
  957. if (gid >= gid_max) return;
  958.  
  959. /**
  960. * main
  961. */
  962.  
  963. m03710s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, l_bin2asc);
  964. }
  965.  
  966. __kernel void m03710_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  967. {
  968. /**
  969. * base
  970. */
  971.  
  972. const u32 gid = get_global_id (0);
  973. const u32 lid = get_local_id (0);
  974. const u32 lsz = get_local_size (0);
  975.  
  976. /**
  977. * bin2asc table
  978. */
  979.  
  980. __local u32 l_bin2asc[256];
  981.  
  982. for (u32 i = lid; i < 256; i += lsz)
  983. {
  984. const u32 i0 = (i >> 0) & 15;
  985. const u32 i1 = (i >> 4) & 15;
  986.  
  987. l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
  988. | ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
  989. }
  990.  
  991. barrier (CLK_LOCAL_MEM_FENCE);
  992.  
  993. if (gid >= gid_max) return;
  994.  
  995. /**
  996. * modifier
  997. */
  998.  
  999. u32 w0[4];
  1000.  
  1001. w0[0] = pws[gid].i[ 0];
  1002. w0[1] = pws[gid].i[ 1];
  1003. w0[2] = pws[gid].i[ 2];
  1004. w0[3] = pws[gid].i[ 3];
  1005.  
  1006. u32 w1[4];
  1007.  
  1008. w1[0] = pws[gid].i[ 4];
  1009. w1[1] = pws[gid].i[ 5];
  1010. w1[2] = pws[gid].i[ 6];
  1011. w1[3] = pws[gid].i[ 7];
  1012.  
  1013. u32 w2[4];
  1014.  
  1015. w2[0] = pws[gid].i[ 8];
  1016. w2[1] = pws[gid].i[ 9];
  1017. w2[2] = pws[gid].i[10];
  1018. w2[3] = pws[gid].i[11];
  1019.  
  1020. u32 w3[4];
  1021.  
  1022. w3[0] = pws[gid].i[12];
  1023. w3[1] = pws[gid].i[13];
  1024. w3[2] = pws[gid].i[14];
  1025. w3[3] = pws[gid].i[15];
  1026.  
  1027. const u32 pw_len = pws[gid].pw_len;
  1028.  
  1029. /**
  1030. * main
  1031. */
  1032.  
  1033. m03710s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, l_bin2asc);
  1034. }
RAW Paste Data