mirror of
https://github.com/hashcat/hashcat
synced 2024-07-20 16:24:12 +02:00
a673aee037
This is a test commit using buffers large enough to handle both passwords and salts up to length 256. It requires changes to the kernel code, which is not included in here. It also requires some of the host code to be modified. Before we're going to modify kernel code to support the larger lengths I want to be sure of: 1. Host code modification is ok (no overflows or underflows) 2. Passwords and Salts are printed correctly to status, outfile, show, left, etc. 3. Performance does not change (or only very minimal) This is not a patch that supports actual cracking both passwords and salts up to length 256, but it can not fail anyway. If if it does, there's no reason to continue to add support for both passwords and salts up to length 256.
48 lines
1.1 KiB
Common Lisp
48 lines
1.1 KiB
Common Lisp
/**
|
|
* Author......: See docs/credits.txt
|
|
* License.....: MIT
|
|
*/
|
|
|
|
#include "inc_hash_constants.h"
|
|
#include "inc_vendor.cl"
|
|
#include "inc_types.cl"
|
|
|
|
#include "inc_rp.h"
|
|
#include "inc_rp.cl"
|
|
|
|
__kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
|
|
{
|
|
const u32 gid = get_global_id (0);
|
|
|
|
if (gid >= gid_max) return;
|
|
|
|
if (rules_buf[0].cmds[0] == RULE_OP_MANGLE_NOOP && rules_buf[0].cmds[1] == 0) return;
|
|
|
|
const u32 pw_len = pws[gid].pw_len;
|
|
|
|
u32 w0[4];
|
|
u32 w1[4];
|
|
|
|
w0[0] = pws[gid].i[ 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);
|
|
|
|
pws_amp[gid].i[0] = w0[0];
|
|
pws_amp[gid].i[1] = w0[1];
|
|
pws_amp[gid].i[2] = w0[2];
|
|
pws_amp[gid].i[3] = w0[3];
|
|
pws_amp[gid].i[4] = w1[0];
|
|
pws_amp[gid].i[5] = w1[1];
|
|
pws_amp[gid].i[6] = w1[2];
|
|
pws_amp[gid].i[7] = w1[3];
|
|
|
|
pws_amp[gid].pw_len = out_len;
|
|
}
|