2 * Author......: Jens Steube <jens.steube@gmail.com>
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"
14 #define COMPARE_S "inc_comp_single.cl"
15 #define COMPARE_M "inc_comp_multi.cl"
24 __constant entry_t pc[1024] =
1052 void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 offset)
1057 w0[0] = w0[0] | append[0] << 8;
1058 w0[1] = append[0] >> 24 | append[1] << 8;
1059 w0[2] = append[1] >> 24 | append[2] << 8;
1060 w0[3] = append[2] >> 24 | append[3] << 8;
1064 w0[0] = w0[0] | append[0] << 16;
1065 w0[1] = append[0] >> 16 | append[1] << 16;
1066 w0[2] = append[1] >> 16 | append[2] << 16;
1067 w0[3] = append[2] >> 16 | append[3] << 16;
1071 w0[0] = w0[0] | append[0] << 24;
1072 w0[1] = append[0] >> 8 | append[1] << 24;
1073 w0[2] = append[1] >> 8 | append[2] << 24;
1074 w0[3] = append[2] >> 8 | append[3] << 24;
1086 void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], const u32 offset)
1095 #if defined IS_AMD || defined IS_GENERIC
1097 const int offset_minus_4 = 4 - (offset & 3);
1099 tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
1100 tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
1101 tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
1102 tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
1103 tmp4 = amd_bytealign (append[4], append[3], offset_minus_4);
1104 tmp5 = amd_bytealign ( 0, append[4], offset_minus_4);
1106 const u32 mod = offset & 3;
1122 const int offset_minus_4 = 4 - (offset & 3);
1124 const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
1126 tmp0 = __byte_perm ( 0, append[0], selector);
1127 tmp1 = __byte_perm (append[0], append[1], selector);
1128 tmp2 = __byte_perm (append[1], append[2], selector);
1129 tmp3 = __byte_perm (append[2], append[3], selector);
1130 tmp4 = __byte_perm (append[3], append[4], selector);
1131 tmp5 = __byte_perm (append[4], 0, selector);
1135 const u32 div = offset / 4;
1139 case 0: w0[0] |= tmp0;
1146 case 1: w0[1] |= tmp0;
1153 case 2: w0[2] |= tmp0;
1160 case 3: w0[3] |= tmp0;
1167 case 4: w1[0] |= tmp0;
1177 void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
1205 SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
1206 SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
1207 SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
1208 SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
1209 SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
1210 SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
1211 SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
1212 SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
1213 SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
1214 SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
1215 SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
1216 SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
1217 SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
1218 SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
1219 SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
1220 SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
1221 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
1222 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
1223 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
1224 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
1229 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
1230 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
1231 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
1232 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
1233 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
1234 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
1235 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
1236 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
1237 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
1238 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
1239 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
1240 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
1241 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
1242 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
1243 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
1244 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
1245 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
1246 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
1247 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
1248 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
1253 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
1254 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
1255 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
1256 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
1257 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
1258 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
1259 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
1260 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
1261 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
1262 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
1263 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
1264 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
1265 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
1266 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
1267 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
1268 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
1269 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
1270 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
1271 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
1272 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
1277 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
1278 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
1279 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
1280 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
1281 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
1282 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
1283 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
1284 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
1285 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
1286 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
1287 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
1288 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
1289 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
1290 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
1291 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
1292 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
1293 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
1294 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
1295 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
1296 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
1305 __kernel void m05800_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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)
1311 const u32 gid = get_global_id (0);
1313 if (gid >= gid_max) return;
1317 word_buf[0] = pws[gid].i[ 0];
1318 word_buf[1] = pws[gid].i[ 1];
1319 word_buf[2] = pws[gid].i[ 2];
1320 word_buf[3] = pws[gid].i[ 3];
1322 const u32 pw_len = pws[gid].pw_len;
1328 u32 salt_len = salt_bufs[salt_pos].salt_len;
1332 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1333 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1334 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1335 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1336 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1342 const u32 pc_len = 1;
1343 const u32 pc_dec = 0x30;
1345 u32 data0[4] = { 0, 0, 0, 0 };
1346 u32 data1[4] = { 0, 0, 0, 0 };
1347 u32 data2[4] = { 0, 0, 0, 0 };
1351 append_word (data0, data1, word_buf, pc_len);
1353 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1360 w0[0] = swap32 (data0[0]);
1361 w0[1] = swap32 (data0[1]);
1362 w0[2] = swap32 (data0[2]);
1363 w0[3] = swap32 (data0[3]);
1364 w1[0] = swap32 (data1[0]);
1365 w1[1] = swap32 (data1[1]);
1366 w1[2] = swap32 (data1[2]);
1367 w1[3] = swap32 (data1[3]);
1368 w2[0] = swap32 (data2[0]);
1369 w2[1] = swap32 (data2[1]);
1375 w3[3] = (pc_len + pw_len + salt_len) * 8;
1379 digest[0] = SHA1M_A;
1380 digest[1] = SHA1M_B;
1381 digest[2] = SHA1M_C;
1382 digest[3] = SHA1M_D;
1383 digest[4] = SHA1M_E;
1385 sha1_transform (w0, w1, w2, w3, digest);
1387 tmps[gid].digest_buf[0] = digest[0];
1388 tmps[gid].digest_buf[1] = digest[1];
1389 tmps[gid].digest_buf[2] = digest[2];
1390 tmps[gid].digest_buf[3] = digest[3];
1391 tmps[gid].digest_buf[4] = digest[4];
1394 __kernel void m05800_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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)
1400 const u32 gid = get_global_id (0);
1401 const u32 lid = get_local_id (0);
1402 const u32 lsz = get_local_size (0);
1405 * cache precomputed conversion table in shared memory
1408 __local entry_t s_pc[1024];
1410 for (u32 i = lid; i < 1024; i += lsz)
1415 barrier (CLK_LOCAL_MEM_FENCE);
1417 if (gid >= gid_max) return;
1425 word_buf[0] = pws[gid].i[ 0];
1426 word_buf[1] = pws[gid].i[ 1];
1427 word_buf[2] = pws[gid].i[ 2];
1428 word_buf[3] = pws[gid].i[ 3];
1430 const u32 pw_len = pws[gid].pw_len;
1434 digest[0] = tmps[gid].digest_buf[0];
1435 digest[1] = tmps[gid].digest_buf[1];
1436 digest[2] = tmps[gid].digest_buf[2];
1437 digest[3] = tmps[gid].digest_buf[3];
1438 digest[4] = tmps[gid].digest_buf[4];
1444 u32 salt_len = salt_bufs[salt_pos].salt_len;
1448 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1449 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1450 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1451 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1452 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1458 for (u32 i = 0, j = loop_pos + 1; i < loop_cnt; i++, j++)
1460 const u32 pc_len = s_pc[j].len;
1461 const u32 pc_dec = s_pc[j].dec;
1463 u32 data0[4] = { 0, 0, 0, 0 };
1464 u32 data1[4] = { 0, 0, 0, 0 };
1465 u32 data2[4] = { 0, 0, 0, 0 };
1469 append_word (data0, data1, word_buf, pc_len);
1471 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1483 w1[1] = swap32 (data0[0]);
1484 w1[2] = swap32 (data0[1]);
1485 w1[3] = swap32 (data0[2]);
1486 w2[0] = swap32 (data0[3]);
1487 w2[1] = swap32 (data1[0]);
1488 w2[2] = swap32 (data1[1]);
1489 w2[3] = swap32 (data1[2]);
1490 w3[0] = swap32 (data1[3]);
1491 w3[1] = swap32 (data2[0]);
1493 w3[3] = (20 + pc_len + pw_len + salt_len) * 8;
1495 digest[0] = SHA1M_A;
1496 digest[1] = SHA1M_B;
1497 digest[2] = SHA1M_C;
1498 digest[3] = SHA1M_D;
1499 digest[4] = SHA1M_E;
1501 sha1_transform (w0, w1, w2, w3, digest);
1504 tmps[gid].digest_buf[0] = digest[0];
1505 tmps[gid].digest_buf[1] = digest[1];
1506 tmps[gid].digest_buf[2] = digest[2];
1507 tmps[gid].digest_buf[3] = digest[3];
1508 tmps[gid].digest_buf[4] = digest[4];
1511 __kernel void m05800_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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)
1517 const u32 gid = get_global_id (0);
1519 if (gid >= gid_max) return;
1521 const u32 lid = get_local_id (0);
1527 const u32 r0 = tmps[gid].digest_buf[DGST_R0];
1528 const u32 r1 = tmps[gid].digest_buf[DGST_R1];
1529 const u32 r2 = tmps[gid].digest_buf[DGST_R2];
1530 const u32 r3 = tmps[gid].digest_buf[DGST_R3];
projects / hashcat.git / blob