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

Replace code to use pure kernel rule engine for slow hashes

Browse Source
This commit is contained in:
jsteube 2017-08-11 16:21:19 +02:00
parent 8a0d21360b
commit 07b54c1257
1 changed files with 35 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

57
OpenCL/amp_a0.cl
View File

@ -6,9 +6,8 @@
#include "inc_hash_constants.h" #include "inc_hash_constants.h"
#include "inc_vendor.cl" #include "inc_vendor.cl"
#include "inc_types.cl" #include "inc_types.cl"
#include "inc_rp.h"
#include "inc_rp_optimized.h" #include "inc_rp.cl"
#include "inc_rp_optimized.cl"
__kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max) __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
{ {
@ -18,30 +17,44 @@ __kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const ke
if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0) return; if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0) return;
/**
* base
*/
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
u32 w0[4]; const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w1[4];
w0[0] = pws[gid].i[ 0]; u32 w[64] = { 0 };
w0[1] = pws[gid].i[ 1];
w0[2] = pws[gid].i[ 2];
w0[3] = pws[gid].i[ 3];
w1[0] = pws[gid].i[ 4];
w1[1] = pws[gid].i[ 5];
w1[2] = pws[gid].i[ 6];
w1[3] = pws[gid].i[ 7];
const u32 out_len = apply_rules (rules_buf[0].cmds, w0, w1, pw_len); for (int idx = 0; idx < pw_lenv; idx++)
{
w[idx] = pws[gid].i[idx];
}
pws_amp[gid].i[0] = w0[0]; /**
pws_amp[gid].i[1] = w0[1]; * do work
pws_amp[gid].i[2] = w0[2]; */
pws_amp[gid].i[3] = w0[3];
pws_amp[gid].i[4] = w1[0]; u32 out_buf[64] = { 0 };
pws_amp[gid].i[5] = w1[1];
pws_amp[gid].i[6] = w1[2]; const u32 out_len = apply_rules (rules_buf[0].cmds, w, pw_len, out_buf);
pws_amp[gid].i[7] = w1[3];
/**
* out
*/
const u32 out_lenv = ceil ((float) out_len / 4);
for (int idx = 0; idx < pw_lenv; idx++)
{
pws_amp[gid].i[idx] = out_buf[idx];
}
for (int idx = pw_lenv; idx < 64; idx++)
{
pws_amp[gid].i[idx] = 0;
}
pws_amp[gid].pw_len = out_len; pws_amp[gid].pw_len = out_len;
} }