2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "inc_hash_constants.h"
9 #include "inc_vendor.cl"
16 #include "inc_hash_functions.cl"
17 #include "inc_types.cl"
18 #include "inc_common.cl"
20 #define COMPARE_S "inc_comp_single.cl"
21 #define COMPARE_M "inc_comp_multi.cl"
30 __constant entry_t pc[1024] =
1058 void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 offset)
1063 w0[0] = w0[0] | append[0] << 8;
1064 w0[1] = append[0] >> 24 | append[1] << 8;
1065 w0[2] = append[1] >> 24 | append[2] << 8;
1066 w0[3] = append[2] >> 24 | append[3] << 8;
1070 w0[0] = w0[0] | append[0] << 16;
1071 w0[1] = append[0] >> 16 | append[1] << 16;
1072 w0[2] = append[1] >> 16 | append[2] << 16;
1073 w0[3] = append[2] >> 16 | append[3] << 16;
1077 w0[0] = w0[0] | append[0] << 24;
1078 w0[1] = append[0] >> 8 | append[1] << 24;
1079 w0[2] = append[1] >> 8 | append[2] << 24;
1080 w0[3] = append[2] >> 8 | append[3] << 24;
1092 void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], const u32 offset)
1101 #if defined IS_AMD || defined IS_GENERIC
1103 const int offset_minus_4 = 4 - (offset & 3);
1105 tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
1106 tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
1107 tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
1108 tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
1109 tmp4 = amd_bytealign (append[4], append[3], offset_minus_4);
1110 tmp5 = amd_bytealign ( 0, append[4], offset_minus_4);
1112 const u32 mod = offset & 3;
1128 const int offset_minus_4 = 4 - (offset & 3);
1130 const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
1132 tmp0 = __byte_perm ( 0, append[0], selector);
1133 tmp1 = __byte_perm (append[0], append[1], selector);
1134 tmp2 = __byte_perm (append[1], append[2], selector);
1135 tmp3 = __byte_perm (append[2], append[3], selector);
1136 tmp4 = __byte_perm (append[3], append[4], selector);
1137 tmp5 = __byte_perm (append[4], 0, selector);
1141 const u32 div = offset / 4;
1145 case 0: w0[0] |= tmp0;
1152 case 1: w0[1] |= tmp0;
1159 case 2: w0[2] |= tmp0;
1166 case 3: w0[3] |= tmp0;
1173 case 4: w1[0] |= tmp0;
1183 void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
1211 SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
1212 SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
1213 SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
1214 SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
1215 SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
1216 SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
1217 SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
1218 SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
1219 SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
1220 SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
1221 SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
1222 SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
1223 SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
1224 SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
1225 SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
1226 SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
1227 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
1228 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
1229 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
1230 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
1235 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
1236 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
1237 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
1238 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
1239 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
1240 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
1241 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
1242 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
1243 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
1244 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
1245 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
1246 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
1247 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
1248 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
1249 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
1250 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
1251 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
1252 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
1253 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
1254 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
1259 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
1260 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
1261 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
1262 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
1263 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
1264 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
1265 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
1266 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
1267 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
1268 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
1269 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
1270 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
1271 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
1272 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
1273 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
1274 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
1275 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
1276 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
1277 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
1278 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
1283 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
1284 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
1285 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
1286 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
1287 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
1288 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
1289 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
1290 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
1291 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
1292 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
1293 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
1294 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
1295 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
1296 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
1297 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
1298 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
1299 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
1300 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
1301 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
1302 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
1311 __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_scryptV_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)
1317 const u32 gid = get_global_id (0);
1319 if (gid >= gid_max) return;
1323 word_buf[0] = pws[gid].i[ 0];
1324 word_buf[1] = pws[gid].i[ 1];
1325 word_buf[2] = pws[gid].i[ 2];
1326 word_buf[3] = pws[gid].i[ 3];
1328 const u32 pw_len = pws[gid].pw_len;
1334 u32 salt_len = salt_bufs[salt_pos].salt_len;
1338 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1339 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1340 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1341 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1342 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1348 const u32 pc_len = 1;
1349 const u32 pc_dec = 0x30;
1351 u32 data0[4] = { 0, 0, 0, 0 };
1352 u32 data1[4] = { 0, 0, 0, 0 };
1353 u32 data2[4] = { 0, 0, 0, 0 };
1357 append_word (data0, data1, word_buf, pc_len);
1359 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1366 w0[0] = swap32 (data0[0]);
1367 w0[1] = swap32 (data0[1]);
1368 w0[2] = swap32 (data0[2]);
1369 w0[3] = swap32 (data0[3]);
1370 w1[0] = swap32 (data1[0]);
1371 w1[1] = swap32 (data1[1]);
1372 w1[2] = swap32 (data1[2]);
1373 w1[3] = swap32 (data1[3]);
1374 w2[0] = swap32 (data2[0]);
1375 w2[1] = swap32 (data2[1]);
1381 w3[3] = (pc_len + pw_len + salt_len) * 8;
1385 digest[0] = SHA1M_A;
1386 digest[1] = SHA1M_B;
1387 digest[2] = SHA1M_C;
1388 digest[3] = SHA1M_D;
1389 digest[4] = SHA1M_E;
1391 sha1_transform (w0, w1, w2, w3, digest);
1393 tmps[gid].digest_buf[0] = digest[0];
1394 tmps[gid].digest_buf[1] = digest[1];
1395 tmps[gid].digest_buf[2] = digest[2];
1396 tmps[gid].digest_buf[3] = digest[3];
1397 tmps[gid].digest_buf[4] = digest[4];
1400 __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_scryptV_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)
1406 const u32 gid = get_global_id (0);
1407 const u32 lid = get_local_id (0);
1408 const u32 lsz = get_local_size (0);
1411 * cache precomputed conversion table in shared memory
1414 __local entry_t s_pc[1024];
1416 for (u32 i = lid; i < 1024; i += lsz)
1421 barrier (CLK_LOCAL_MEM_FENCE);
1423 if (gid >= gid_max) return;
1431 word_buf[0] = pws[gid].i[ 0];
1432 word_buf[1] = pws[gid].i[ 1];
1433 word_buf[2] = pws[gid].i[ 2];
1434 word_buf[3] = pws[gid].i[ 3];
1436 const u32 pw_len = pws[gid].pw_len;
1440 digest[0] = tmps[gid].digest_buf[0];
1441 digest[1] = tmps[gid].digest_buf[1];
1442 digest[2] = tmps[gid].digest_buf[2];
1443 digest[3] = tmps[gid].digest_buf[3];
1444 digest[4] = tmps[gid].digest_buf[4];
1450 u32 salt_len = salt_bufs[salt_pos].salt_len;
1454 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1455 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1456 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1457 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1458 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1464 for (u32 i = 0, j = loop_pos + 1; i < loop_cnt; i++, j++)
1466 const u32 pc_len = s_pc[j].len;
1467 const u32 pc_dec = s_pc[j].dec;
1469 u32 data0[4] = { 0, 0, 0, 0 };
1470 u32 data1[4] = { 0, 0, 0, 0 };
1471 u32 data2[4] = { 0, 0, 0, 0 };
1475 append_word (data0, data1, word_buf, pc_len);
1477 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1489 w1[1] = swap32 (data0[0]);
1490 w1[2] = swap32 (data0[1]);
1491 w1[3] = swap32 (data0[2]);
1492 w2[0] = swap32 (data0[3]);
1493 w2[1] = swap32 (data1[0]);
1494 w2[2] = swap32 (data1[1]);
1495 w2[3] = swap32 (data1[2]);
1496 w3[0] = swap32 (data1[3]);
1497 w3[1] = swap32 (data2[0]);
1499 w3[3] = (20 + pc_len + pw_len + salt_len) * 8;
1501 digest[0] = SHA1M_A;
1502 digest[1] = SHA1M_B;
1503 digest[2] = SHA1M_C;
1504 digest[3] = SHA1M_D;
1505 digest[4] = SHA1M_E;
1507 sha1_transform (w0, w1, w2, w3, digest);
1510 tmps[gid].digest_buf[0] = digest[0];
1511 tmps[gid].digest_buf[1] = digest[1];
1512 tmps[gid].digest_buf[2] = digest[2];
1513 tmps[gid].digest_buf[3] = digest[3];
1514 tmps[gid].digest_buf[4] = digest[4];
1517 __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_scryptV_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)
1523 const u32 gid = get_global_id (0);
1525 if (gid >= gid_max) return;
1527 const u32 lid = get_local_id (0);
1533 const u32 r0 = tmps[gid].digest_buf[DGST_R0];
1534 const u32 r1 = tmps[gid].digest_buf[DGST_R1];
1535 const u32 r2 = tmps[gid].digest_buf[DGST_R2];
1536 const u32 r3 = tmps[gid].digest_buf[DGST_R3];
projects / hashcat.git / blob