arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2025-01-11 00:01:16 +00:00
Code Issues Releases Wiki Activity

A bit a different _comp kernel iteration for WPA which can have lots of handshakes

Browse Source
This commit is contained in:
jsteube 2017-03-10 19:54:00 +01:00
parent c304f3163f
commit 2bc65c2c4d
2 changed files with 210 additions and 196 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
OpenCL/m02500.cl
View File

@ -854,8 +854,8 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
hmac_sha1_pad_S (w0, w1, w2, w3, ipad_l0, opad_l0);
for (u32 digest_pos = 0; digest_pos < digests_cnt; digest_pos++)
{
const u32 digest_pos = loop_pos;
const u32 digest_cur = digests_offset + digest_pos;
__global wpa_t *wpa = &wpa_bufs[digest_cur];
@ -1096,5 +1096,4 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, digest_cur, gid, 0);
}
}
}
}

15
src/opencl.c
View File

@ -1289,9 +1289,24 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
if (run_comp == true)
{
u32 loops_cnt = 1;
if (hashconfig->hash_mode == 2500)
{
loops_cnt = hashes->salts_buf[salt_pos].digests_cnt;
}
for (u32 loops_pos = 0; loops_pos < loops_cnt; loops_pos++)
{
device_param->kernel_params_buf32[28] = loops_pos;
device_param->kernel_params_buf32[29] = loops_cnt;
CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_cnt, false, 0);
if (CL_rc == -1) return -1;
if (status_ctx->run_thread_level2 == false) break;
}
}
}