2 * Author......: Jens Steube <jens.steube@gmail.com>
8 #include "include/constants.h"
9 #include "include/kernel_vendor.h"
16 #include "include/kernel_functions.c"
17 #include "OpenCL/types_ocl.c"
18 #include "OpenCL/common.c"
20 #define COMPARE_S "OpenCL/check_single_comp4.c"
21 #define COMPARE_M "OpenCL/check_multi_comp4.c"
30 __constant entry_t pc[1024] =
1058 static 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 static 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 static 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 __attribute__((reqd_work_group_size (64, 1, 1))) 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 rules_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 __attribute__((reqd_work_group_size (64, 1, 1))) 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 rules_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);
1410 word_buf[0] = pws[gid].i[ 0];
1411 word_buf[1] = pws[gid].i[ 1];
1412 word_buf[2] = pws[gid].i[ 2];
1413 word_buf[3] = pws[gid].i[ 3];
1415 const u32 pw_len = pws[gid].pw_len;
1417 const u32 lid = get_local_id (0);
1421 digest[0] = tmps[gid].digest_buf[0];
1422 digest[1] = tmps[gid].digest_buf[1];
1423 digest[2] = tmps[gid].digest_buf[2];
1424 digest[3] = tmps[gid].digest_buf[3];
1425 digest[4] = tmps[gid].digest_buf[4];
1428 * cache precomputed conversion table in shared memory
1431 __local entry_t s_pc[1024];
1433 const u32 lid16 = lid * 16;
1435 s_pc[lid16 + 0] = pc[lid16 + 0];
1436 s_pc[lid16 + 1] = pc[lid16 + 1];
1437 s_pc[lid16 + 2] = pc[lid16 + 2];
1438 s_pc[lid16 + 3] = pc[lid16 + 3];
1439 s_pc[lid16 + 4] = pc[lid16 + 4];
1440 s_pc[lid16 + 5] = pc[lid16 + 5];
1441 s_pc[lid16 + 6] = pc[lid16 + 6];
1442 s_pc[lid16 + 7] = pc[lid16 + 7];
1443 s_pc[lid16 + 8] = pc[lid16 + 8];
1444 s_pc[lid16 + 9] = pc[lid16 + 9];
1445 s_pc[lid16 + 10] = pc[lid16 + 10];
1446 s_pc[lid16 + 11] = pc[lid16 + 11];
1447 s_pc[lid16 + 12] = pc[lid16 + 12];
1448 s_pc[lid16 + 13] = pc[lid16 + 13];
1449 s_pc[lid16 + 14] = pc[lid16 + 14];
1450 s_pc[lid16 + 15] = pc[lid16 + 15];
1452 barrier (CLK_LOCAL_MEM_FENCE);
1454 if (gid >= gid_max) return;
1460 u32 salt_len = salt_bufs[salt_pos].salt_len;
1464 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1465 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1466 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1467 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1468 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1474 for (u32 i = 0, j = loop_pos + 1; i < loop_cnt; i++, j++)
1476 const u32 pc_len = s_pc[j].len;
1477 const u32 pc_dec = s_pc[j].dec;
1479 u32 data0[4] = { 0, 0, 0, 0 };
1480 u32 data1[4] = { 0, 0, 0, 0 };
1481 u32 data2[4] = { 0, 0, 0, 0 };
1485 append_word (data0, data1, word_buf, pc_len);
1487 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1499 w1[1] = swap32 (data0[0]);
1500 w1[2] = swap32 (data0[1]);
1501 w1[3] = swap32 (data0[2]);
1502 w2[0] = swap32 (data0[3]);
1503 w2[1] = swap32 (data1[0]);
1504 w2[2] = swap32 (data1[1]);
1505 w2[3] = swap32 (data1[2]);
1506 w3[0] = swap32 (data1[3]);
1507 w3[1] = swap32 (data2[0]);
1509 w3[3] = (20 + pc_len + pw_len + salt_len) * 8;
1511 digest[0] = SHA1M_A;
1512 digest[1] = SHA1M_B;
1513 digest[2] = SHA1M_C;
1514 digest[3] = SHA1M_D;
1515 digest[4] = SHA1M_E;
1517 sha1_transform (w0, w1, w2, w3, digest);
1520 tmps[gid].digest_buf[0] = digest[0];
1521 tmps[gid].digest_buf[1] = digest[1];
1522 tmps[gid].digest_buf[2] = digest[2];
1523 tmps[gid].digest_buf[3] = digest[3];
1524 tmps[gid].digest_buf[4] = digest[4];
1527 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1533 const u32 gid = get_global_id (0);
1535 if (gid >= gid_max) return;
1537 const u32 lid = get_local_id (0);
1543 const u32 r0 = tmps[gid].digest_buf[DGST_R0];
1544 const u32 r1 = tmps[gid].digest_buf[DGST_R1];
1545 const u32 r2 = tmps[gid].digest_buf[DGST_R2];
1546 const u32 r3 = tmps[gid].digest_buf[DGST_R3];