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 "types_ocl.c"
20 #define COMPARE_S "check_single_comp4.c"
21 #define COMPARE_M "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)
1097 w0[0] = w0[0] | append[0] << 16;
1098 w0[1] = append[0] >> 16 | append[1] << 16;
1099 w0[2] = append[1] >> 16 | append[2] << 16;
1100 w0[3] = append[2] >> 16 | append[3] << 16;
1101 w1[0] = append[3] >> 16 | append[4] << 16;
1102 w1[1] = append[4] >> 16;
1106 w0[0] = w0[0] | append[0] << 24;
1107 w0[1] = append[0] >> 8 | append[1] << 24;
1108 w0[2] = append[1] >> 8 | append[2] << 24;
1109 w0[3] = append[2] >> 8 | append[3] << 24;
1110 w1[0] = append[3] >> 8 | append[4] << 24;
1111 w1[1] = append[4] >> 8;
1123 w0[1] = w0[1] | append[0] << 8;
1124 w0[2] = append[0] >> 24 | append[1] << 8;
1125 w0[3] = append[1] >> 24 | append[2] << 8;
1126 w1[0] = append[2] >> 24 | append[3] << 8;
1127 w1[1] = append[3] >> 24 | append[4] << 8;
1128 w1[2] = append[4] >> 24;
1132 w0[1] = w0[1] | append[0] << 16;
1133 w0[2] = append[0] >> 16 | append[1] << 16;
1134 w0[3] = append[1] >> 16 | append[2] << 16;
1135 w1[0] = append[2] >> 16 | append[3] << 16;
1136 w1[1] = append[3] >> 16 | append[4] << 16;
1137 w1[2] = append[4] >> 16;
1141 w0[1] = w0[1] | append[0] << 24;
1142 w0[2] = append[0] >> 8 | append[1] << 24;
1143 w0[3] = append[1] >> 8 | append[2] << 24;
1144 w1[0] = append[2] >> 8 | append[3] << 24;
1145 w1[1] = append[3] >> 8 | append[4] << 24;
1146 w1[2] = append[4] >> 8;
1158 w0[2] = w0[2] | append[0] << 8;
1159 w0[3] = append[0] >> 24 | append[1] << 8;
1160 w1[0] = append[1] >> 24 | append[2] << 8;
1161 w1[1] = append[2] >> 24 | append[3] << 8;
1162 w1[2] = append[3] >> 24 | append[4] << 8;
1163 w1[3] = append[4] >> 24;
1167 w0[2] = w0[2] | append[0] << 16;
1168 w0[3] = append[0] >> 16 | append[1] << 16;
1169 w1[0] = append[1] >> 16 | append[2] << 16;
1170 w1[1] = append[2] >> 16 | append[3] << 16;
1171 w1[2] = append[3] >> 16 | append[4] << 16;
1172 w1[3] = append[4] >> 16;
1176 w0[2] = w0[2] | append[0] << 24;
1177 w0[3] = append[0] >> 8 | append[1] << 24;
1178 w1[0] = append[1] >> 8 | append[2] << 24;
1179 w1[1] = append[2] >> 8 | append[3] << 24;
1180 w1[2] = append[3] >> 8 | append[4] << 24;
1181 w1[3] = append[4] >> 8;
1193 w0[3] = w0[3] | append[0] << 8;
1194 w1[0] = append[0] >> 24 | append[1] << 8;
1195 w1[1] = append[1] >> 24 | append[2] << 8;
1196 w1[2] = append[2] >> 24 | append[3] << 8;
1197 w1[3] = append[3] >> 24 | append[4] << 8;
1198 w2[0] = append[4] >> 24;
1202 w0[3] = w0[3] | append[0] << 16;
1203 w1[0] = append[0] >> 16 | append[1] << 16;
1204 w1[1] = append[1] >> 16 | append[2] << 16;
1205 w1[2] = append[2] >> 16 | append[3] << 16;
1206 w1[3] = append[3] >> 16 | append[4] << 16;
1207 w2[0] = append[4] >> 16;
1211 w0[3] = w0[3] | append[0] << 24;
1212 w1[0] = append[0] >> 8 | append[1] << 24;
1213 w1[1] = append[1] >> 8 | append[2] << 24;
1214 w1[2] = append[2] >> 8 | append[3] << 24;
1215 w1[3] = append[3] >> 8 | append[4] << 24;
1216 w2[0] = append[4] >> 8;
1228 w1[0] = w1[0] | append[0] << 8;
1229 w1[1] = append[0] >> 24 | append[1] << 8;
1230 w1[2] = append[1] >> 24 | append[2] << 8;
1231 w1[3] = append[2] >> 24 | append[3] << 8;
1232 w2[0] = append[3] >> 24 | append[4] << 8;
1233 w2[1] = append[4] >> 24;
1237 w1[0] = w1[0] | append[0] << 16;
1238 w1[1] = append[0] >> 16 | append[1] << 16;
1239 w1[2] = append[1] >> 16 | append[2] << 16;
1240 w1[3] = append[2] >> 16 | append[3] << 16;
1241 w2[0] = append[3] >> 16 | append[4] << 16;
1242 w2[1] = append[4] >> 16;
1246 w1[0] = w1[0] | append[0] << 24;
1247 w1[1] = append[0] >> 8 | append[1] << 24;
1248 w1[2] = append[1] >> 8 | append[2] << 24;
1249 w1[3] = append[2] >> 8 | append[3] << 24;
1250 w2[0] = append[3] >> 8 | append[4] << 24;
1251 w2[1] = append[4] >> 8;
1256 static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
1284 SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
1285 SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
1286 SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
1287 SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
1288 SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
1289 SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
1290 SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
1291 SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
1292 SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
1293 SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
1294 SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
1295 SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
1296 SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
1297 SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
1298 SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
1299 SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
1300 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
1301 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
1302 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
1303 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
1308 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
1309 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
1310 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
1311 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
1312 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
1313 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
1314 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
1315 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
1316 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
1317 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
1318 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
1319 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
1320 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
1321 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
1322 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
1323 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
1324 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
1325 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
1326 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
1327 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
1332 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
1333 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
1334 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
1335 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
1336 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
1337 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
1338 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
1339 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
1340 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
1341 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
1342 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
1343 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
1344 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
1345 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
1346 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
1347 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
1348 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
1349 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
1350 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
1351 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
1356 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
1357 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
1358 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
1359 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
1360 w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
1361 w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
1362 w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
1363 w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
1364 w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
1365 w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
1366 w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
1367 w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
1368 w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
1369 w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
1370 wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
1371 wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
1372 wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
1373 wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
1374 we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
1375 wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
1384 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_init (__global pw_t *pws, __global gpu_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)
1390 const u32 gid = get_global_id (0);
1392 if (gid >= gid_max) return;
1396 word_buf[0] = pws[gid].i[ 0];
1397 word_buf[1] = pws[gid].i[ 1];
1398 word_buf[2] = pws[gid].i[ 2];
1399 word_buf[3] = pws[gid].i[ 3];
1401 const u32 pw_len = pws[gid].pw_len;
1407 u32 salt_len = salt_bufs[salt_pos].salt_len;
1411 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1412 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1413 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1414 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1415 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1421 const u32 pc_len = 1;
1422 const u32 pc_dec = 0x30;
1424 u32 data0[4] = { 0, 0, 0, 0 };
1425 u32 data1[4] = { 0, 0, 0, 0 };
1426 u32 data2[4] = { 0, 0, 0, 0 };
1430 append_word (data0, data1, word_buf, pc_len);
1432 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1439 w0[0] = swap32 (data0[0]);
1440 w0[1] = swap32 (data0[1]);
1441 w0[2] = swap32 (data0[2]);
1442 w0[3] = swap32 (data0[3]);
1443 w1[0] = swap32 (data1[0]);
1444 w1[1] = swap32 (data1[1]);
1445 w1[2] = swap32 (data1[2]);
1446 w1[3] = swap32 (data1[3]);
1447 w2[0] = swap32 (data2[0]);
1448 w2[1] = swap32 (data2[1]);
1454 w3[3] = (pc_len + pw_len + salt_len) * 8;
1458 digest[0] = SHA1M_A;
1459 digest[1] = SHA1M_B;
1460 digest[2] = SHA1M_C;
1461 digest[3] = SHA1M_D;
1462 digest[4] = SHA1M_E;
1464 sha1_transform (w0, w1, w2, w3, digest);
1466 tmps[gid].digest_buf[0] = digest[0];
1467 tmps[gid].digest_buf[1] = digest[1];
1468 tmps[gid].digest_buf[2] = digest[2];
1469 tmps[gid].digest_buf[3] = digest[3];
1470 tmps[gid].digest_buf[4] = digest[4];
1473 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_loop (__global pw_t *pws, __global gpu_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)
1479 const u32 gid = get_global_id (0);
1483 word_buf[0] = pws[gid].i[ 0];
1484 word_buf[1] = pws[gid].i[ 1];
1485 word_buf[2] = pws[gid].i[ 2];
1486 word_buf[3] = pws[gid].i[ 3];
1488 const u32 pw_len = pws[gid].pw_len;
1490 const u32 lid = get_local_id (0);
1494 digest[0] = tmps[gid].digest_buf[0];
1495 digest[1] = tmps[gid].digest_buf[1];
1496 digest[2] = tmps[gid].digest_buf[2];
1497 digest[3] = tmps[gid].digest_buf[3];
1498 digest[4] = tmps[gid].digest_buf[4];
1501 * cache precomputed conversion table in shared memory
1504 __local entry_t s_pc[1024];
1506 const u32 lid16 = lid * 16;
1508 s_pc[lid16 + 0] = pc[lid16 + 0];
1509 s_pc[lid16 + 1] = pc[lid16 + 1];
1510 s_pc[lid16 + 2] = pc[lid16 + 2];
1511 s_pc[lid16 + 3] = pc[lid16 + 3];
1512 s_pc[lid16 + 4] = pc[lid16 + 4];
1513 s_pc[lid16 + 5] = pc[lid16 + 5];
1514 s_pc[lid16 + 6] = pc[lid16 + 6];
1515 s_pc[lid16 + 7] = pc[lid16 + 7];
1516 s_pc[lid16 + 8] = pc[lid16 + 8];
1517 s_pc[lid16 + 9] = pc[lid16 + 9];
1518 s_pc[lid16 + 10] = pc[lid16 + 10];
1519 s_pc[lid16 + 11] = pc[lid16 + 11];
1520 s_pc[lid16 + 12] = pc[lid16 + 12];
1521 s_pc[lid16 + 13] = pc[lid16 + 13];
1522 s_pc[lid16 + 14] = pc[lid16 + 14];
1523 s_pc[lid16 + 15] = pc[lid16 + 15];
1525 barrier (CLK_LOCAL_MEM_FENCE);
1527 if (gid >= gid_max) return;
1533 u32 salt_len = salt_bufs[salt_pos].salt_len;
1537 salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
1538 salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
1539 salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
1540 salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
1541 salt_buf[4] = salt_bufs[salt_pos].salt_buf[4];
1547 for (u32 i = 0, j = loop_pos + 1; i < loop_cnt; i++, j++)
1549 const u32 pc_len = s_pc[j].len;
1550 const u32 pc_dec = s_pc[j].dec;
1552 u32 data0[4] = { 0, 0, 0, 0 };
1553 u32 data1[4] = { 0, 0, 0, 0 };
1554 u32 data2[4] = { 0, 0, 0, 0 };
1558 append_word (data0, data1, word_buf, pc_len);
1560 append_salt (data0, data1, data2, salt_buf, pc_len + pw_len);
1572 w1[1] = swap32 (data0[0]);
1573 w1[2] = swap32 (data0[1]);
1574 w1[3] = swap32 (data0[2]);
1575 w2[0] = swap32 (data0[3]);
1576 w2[1] = swap32 (data1[0]);
1577 w2[2] = swap32 (data1[1]);
1578 w2[3] = swap32 (data1[2]);
1579 w3[0] = swap32 (data1[3]);
1580 w3[1] = swap32 (data2[0]);
1582 w3[3] = (20 + pc_len + pw_len + salt_len) * 8;
1584 digest[0] = SHA1M_A;
1585 digest[1] = SHA1M_B;
1586 digest[2] = SHA1M_C;
1587 digest[3] = SHA1M_D;
1588 digest[4] = SHA1M_E;
1590 sha1_transform (w0, w1, w2, w3, digest);
1593 tmps[gid].digest_buf[0] = digest[0];
1594 tmps[gid].digest_buf[1] = digest[1];
1595 tmps[gid].digest_buf[2] = digest[2];
1596 tmps[gid].digest_buf[3] = digest[3];
1597 tmps[gid].digest_buf[4] = digest[4];
1600 __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_comp (__global pw_t *pws, __global gpu_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)
1606 const u32 gid = get_global_id (0);
1608 if (gid >= gid_max) return;
1610 const u32 lid = get_local_id (0);
1616 const u32 r0 = tmps[gid].digest_buf[DGST_R0];
1617 const u32 r1 = tmps[gid].digest_buf[DGST_R1];
1618 const u32 r2 = tmps[gid].digest_buf[DGST_R2];
1619 const u32 r3 = tmps[gid].digest_buf[DGST_R3];