arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-22 22:58:30 +00:00
Code Issues Releases Wiki Activity

Cleanup -m 1460 kernels to latest standard

Browse Source
This commit is contained in:
Jens Steube 2016-03-03 17:10:43 +01:00
parent 4f1d33216b
commit 180f71f291
2 changed files with 160 additions and 148 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

172
OpenCL/m01460_a0.cl
View File

@ -246,18 +246,16 @@ __kernel void m01460_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
if (gid >= gid_max) return; if (gid >= gid_max) return;
u32 pw_buf0[4]; u32 pw_buf0[4];
pw_buf0[0] = pws[gid].i[ 0];
pw_buf0[1] = pws[gid].i[ 1];
pw_buf0[2] = pws[gid].i[ 2];
pw_buf0[3] = pws[gid].i[ 3];
u32 pw_buf1[4]; u32 pw_buf1[4];
pw_buf1[0] = pws[gid].i[ 4]; pw_buf0[0] = pws[gid].i[0];
pw_buf1[1] = pws[gid].i[ 5]; pw_buf0[1] = pws[gid].i[1];
pw_buf1[2] = pws[gid].i[ 6]; pw_buf0[2] = pws[gid].i[2];
pw_buf1[3] = pws[gid].i[ 7]; pw_buf0[3] = pws[gid].i[3];
pw_buf1[0] = pws[gid].i[4];
pw_buf1[1] = pws[gid].i[5];
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
@ -266,50 +264,54 @@ __kernel void m01460_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
*/ */
u32 salt_buf0[4]; u32 salt_buf0[4];
salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
u32 salt_buf1[4]; u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[4];
salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4]; salt_buf0[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 0]);
salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5]; salt_buf0[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 1]);
salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6]; salt_buf0[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 2]);
salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7]; salt_buf0[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 3]);
salt_buf1[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 4]);
salt_buf1[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 5]);
salt_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 6]);
salt_buf1[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 7]);
salt_buf2[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 8]);
salt_buf2[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 9]);
salt_buf2[2] = swap32_S (salt_bufs[salt_pos].salt_buf[10]);
salt_buf2[3] = swap32_S (salt_bufs[salt_pos].salt_buf[11]);
salt_buf3[0] = swap32_S (salt_bufs[salt_pos].salt_buf[12]);
salt_buf3[1] = swap32_S (salt_bufs[salt_pos].salt_buf[13]);
salt_buf3[2] = swap32_S (salt_bufs[salt_pos].salt_buf[14]);
salt_buf3[3] = swap32_S (salt_bufs[salt_pos].salt_buf[15]);
const u32 salt_len = salt_bufs[salt_pos].salt_len;
/** /**
* pads * pads
*/ */
u32x w0_t[4]; u32x w0_t[4];
w0_t[0] = swap32 (salt_buf0[0]);
w0_t[1] = swap32 (salt_buf0[1]);
w0_t[2] = swap32 (salt_buf0[2]);
w0_t[3] = swap32 (salt_buf0[3]);
u32x w1_t[4]; u32x w1_t[4];
w1_t[0] = swap32 (salt_buf1[0]);
w1_t[1] = swap32 (salt_buf1[1]);
w1_t[2] = swap32 (salt_buf1[2]);
w1_t[3] = swap32 (salt_buf1[3]);
u32x w2_t[4]; u32x w2_t[4];
w2_t[0] = 0;
w2_t[1] = 0;
w2_t[2] = 0;
w2_t[3] = 0;
u32x w3_t[4]; u32x w3_t[4];
w3_t[0] = 0; w0_t[0] = salt_buf0[0];
w3_t[1] = 0; w0_t[1] = salt_buf0[1];
w3_t[2] = 0; w0_t[2] = salt_buf0[2];
w3_t[3] = 0; w0_t[3] = salt_buf0[3];
w1_t[0] = salt_buf1[0];
w1_t[1] = salt_buf1[1];
w1_t[2] = salt_buf1[2];
w1_t[3] = salt_buf1[3];
w2_t[0] = salt_buf2[0];
w2_t[1] = salt_buf2[1];
w2_t[2] = salt_buf2[2];
w2_t[3] = salt_buf2[3];
w3_t[0] = salt_buf3[0];
w3_t[1] = salt_buf3[1];
w3_t[2] = salt_buf3[2];
w3_t[3] = salt_buf3[3];
u32x ipad[8]; u32x ipad[8];
u32x opad[8]; u32x opad[8];
@ -381,18 +383,16 @@ __kernel void m01460_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
if (gid >= gid_max) return; if (gid >= gid_max) return;
u32 pw_buf0[4]; u32 pw_buf0[4];
pw_buf0[0] = pws[gid].i[ 0];
pw_buf0[1] = pws[gid].i[ 1];
pw_buf0[2] = pws[gid].i[ 2];
pw_buf0[3] = pws[gid].i[ 3];
u32 pw_buf1[4]; u32 pw_buf1[4];
pw_buf1[0] = pws[gid].i[ 4]; pw_buf0[0] = pws[gid].i[0];
pw_buf1[1] = pws[gid].i[ 5]; pw_buf0[1] = pws[gid].i[1];
pw_buf1[2] = pws[gid].i[ 6]; pw_buf0[2] = pws[gid].i[2];
pw_buf1[3] = pws[gid].i[ 7]; pw_buf0[3] = pws[gid].i[3];
pw_buf1[0] = pws[gid].i[4];
pw_buf1[1] = pws[gid].i[5];
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len; const u32 pw_len = pws[gid].pw_len;
@ -401,50 +401,54 @@ __kernel void m01460_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
*/ */
u32 salt_buf0[4]; u32 salt_buf0[4];
salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
u32 salt_buf1[4]; u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[4];
salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4]; salt_buf0[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 0]);
salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5]; salt_buf0[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 1]);
salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6]; salt_buf0[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 2]);
salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7]; salt_buf0[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 3]);
salt_buf1[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 4]);
salt_buf1[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 5]);
salt_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 6]);
salt_buf1[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 7]);
salt_buf2[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 8]);
salt_buf2[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 9]);
salt_buf2[2] = swap32_S (salt_bufs[salt_pos].salt_buf[10]);
salt_buf2[3] = swap32_S (salt_bufs[salt_pos].salt_buf[11]);
salt_buf3[0] = swap32_S (salt_bufs[salt_pos].salt_buf[12]);
salt_buf3[1] = swap32_S (salt_bufs[salt_pos].salt_buf[13]);
salt_buf3[2] = swap32_S (salt_bufs[salt_pos].salt_buf[14]);
salt_buf3[3] = swap32_S (salt_bufs[salt_pos].salt_buf[15]);
const u32 salt_len = salt_bufs[salt_pos].salt_len;
/** /**
* pads * pads
*/ */
u32x w0_t[4]; u32x w0_t[4];
w0_t[0] = swap32 (salt_buf0[0]);
w0_t[1] = swap32 (salt_buf0[1]);
w0_t[2] = swap32 (salt_buf0[2]);
w0_t[3] = swap32 (salt_buf0[3]);
u32x w1_t[4]; u32x w1_t[4];
w1_t[0] = swap32 (salt_buf1[0]);
w1_t[1] = swap32 (salt_buf1[1]);
w1_t[2] = swap32 (salt_buf1[2]);
w1_t[3] = swap32 (salt_buf1[3]);
u32x w2_t[4]; u32x w2_t[4];
w2_t[0] = 0;
w2_t[1] = 0;
w2_t[2] = 0;
w2_t[3] = 0;
u32x w3_t[4]; u32x w3_t[4];
w3_t[0] = 0; w0_t[0] = salt_buf0[0];
w3_t[1] = 0; w0_t[1] = salt_buf0[1];
w3_t[2] = 0; w0_t[2] = salt_buf0[2];
w3_t[3] = 0; w0_t[3] = salt_buf0[3];
w1_t[0] = salt_buf1[0];
w1_t[1] = salt_buf1[1];
w1_t[2] = salt_buf1[2];
w1_t[3] = salt_buf1[3];
w2_t[0] = salt_buf2[0];
w2_t[1] = salt_buf2[1];
w2_t[2] = salt_buf2[2];
w2_t[3] = salt_buf2[3];
w3_t[0] = salt_buf3[0];
w3_t[1] = salt_buf3[1];
w3_t[2] = salt_buf3[2];
w3_t[3] = salt_buf3[3];
u32x ipad[8]; u32x ipad[8];
u32x opad[8]; u32x opad[8];

136
OpenCL/m01460_a3.cl
View File

@ -241,50 +241,54 @@ static void m01460m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
*/ */
u32 salt_buf0[4]; u32 salt_buf0[4];
salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
u32 salt_buf1[4]; u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[4];
salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4]; salt_buf0[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 0]);
salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5]; salt_buf0[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 1]);
salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6]; salt_buf0[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 2]);
salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7]; salt_buf0[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 3]);
salt_buf1[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 4]);
salt_buf1[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 5]);
salt_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 6]);
salt_buf1[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 7]);
salt_buf2[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 8]);
salt_buf2[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 9]);
salt_buf2[2] = swap32_S (salt_bufs[salt_pos].salt_buf[10]);
salt_buf2[3] = swap32_S (salt_bufs[salt_pos].salt_buf[11]);
salt_buf3[0] = swap32_S (salt_bufs[salt_pos].salt_buf[12]);
salt_buf3[1] = swap32_S (salt_bufs[salt_pos].salt_buf[13]);
salt_buf3[2] = swap32_S (salt_bufs[salt_pos].salt_buf[14]);
salt_buf3[3] = swap32_S (salt_bufs[salt_pos].salt_buf[15]);
const u32 salt_len = salt_bufs[salt_pos].salt_len;
/** /**
* pads * pads
*/ */
u32x w0_t[4]; u32x w0_t[4];
w0_t[0] = swap32 (salt_buf0[0]);
w0_t[1] = swap32 (salt_buf0[1]);
w0_t[2] = swap32 (salt_buf0[2]);
w0_t[3] = swap32 (salt_buf0[3]);
u32x w1_t[4]; u32x w1_t[4];
w1_t[0] = swap32 (salt_buf1[0]);
w1_t[1] = swap32 (salt_buf1[1]);
w1_t[2] = swap32 (salt_buf1[2]);
w1_t[3] = swap32 (salt_buf1[3]);
u32x w2_t[4]; u32x w2_t[4];
w2_t[0] = 0;
w2_t[1] = 0;
w2_t[2] = 0;
w2_t[3] = 0;
u32x w3_t[4]; u32x w3_t[4];
w3_t[0] = 0; w0_t[0] = salt_buf0[0];
w3_t[1] = 0; w0_t[1] = salt_buf0[1];
w3_t[2] = 0; w0_t[2] = salt_buf0[2];
w3_t[3] = 0; w0_t[3] = salt_buf0[3];
w1_t[0] = salt_buf1[0];
w1_t[1] = salt_buf1[1];
w1_t[2] = salt_buf1[2];
w1_t[3] = salt_buf1[3];
w2_t[0] = salt_buf2[0];
w2_t[1] = salt_buf2[1];
w2_t[2] = salt_buf2[2];
w2_t[3] = salt_buf2[3];
w3_t[0] = salt_buf3[0];
w3_t[1] = salt_buf3[1];
w3_t[2] = salt_buf3[2];
w3_t[3] = salt_buf3[3];
u32x ipad[8]; u32x ipad[8];
u32x opad[8]; u32x opad[8];
@ -342,50 +346,54 @@ static void m01460s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_le
*/ */
u32 salt_buf0[4]; u32 salt_buf0[4];
salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
u32 salt_buf1[4]; u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[4];
salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4]; salt_buf0[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 0]);
salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5]; salt_buf0[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 1]);
salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6]; salt_buf0[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 2]);
salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7]; salt_buf0[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 3]);
salt_buf1[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 4]);
salt_buf1[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 5]);
salt_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf[ 6]);
salt_buf1[3] = swap32_S (salt_bufs[salt_pos].salt_buf[ 7]);
salt_buf2[0] = swap32_S (salt_bufs[salt_pos].salt_buf[ 8]);
salt_buf2[1] = swap32_S (salt_bufs[salt_pos].salt_buf[ 9]);
salt_buf2[2] = swap32_S (salt_bufs[salt_pos].salt_buf[10]);
salt_buf2[3] = swap32_S (salt_bufs[salt_pos].salt_buf[11]);
salt_buf3[0] = swap32_S (salt_bufs[salt_pos].salt_buf[12]);
salt_buf3[1] = swap32_S (salt_bufs[salt_pos].salt_buf[13]);
salt_buf3[2] = swap32_S (salt_bufs[salt_pos].salt_buf[14]);
salt_buf3[3] = swap32_S (salt_bufs[salt_pos].salt_buf[15]);
const u32 salt_len = salt_bufs[salt_pos].salt_len;
/** /**
* pads * pads
*/ */
u32x w0_t[4]; u32x w0_t[4];
w0_t[0] = swap32 (salt_buf0[0]);
w0_t[1] = swap32 (salt_buf0[1]);
w0_t[2] = swap32 (salt_buf0[2]);
w0_t[3] = swap32 (salt_buf0[3]);
u32x w1_t[4]; u32x w1_t[4];
w1_t[0] = swap32 (salt_buf1[0]);
w1_t[1] = swap32 (salt_buf1[1]);
w1_t[2] = swap32 (salt_buf1[2]);
w1_t[3] = swap32 (salt_buf1[3]);
u32x w2_t[4]; u32x w2_t[4];
w2_t[0] = 0;
w2_t[1] = 0;
w2_t[2] = 0;
w2_t[3] = 0;
u32x w3_t[4]; u32x w3_t[4];
w3_t[0] = 0; w0_t[0] = salt_buf0[0];
w3_t[1] = 0; w0_t[1] = salt_buf0[1];
w3_t[2] = 0; w0_t[2] = salt_buf0[2];
w3_t[3] = 0; w0_t[3] = salt_buf0[3];
w1_t[0] = salt_buf1[0];
w1_t[1] = salt_buf1[1];
w1_t[2] = salt_buf1[2];
w1_t[3] = salt_buf1[3];
w2_t[0] = salt_buf2[0];
w2_t[1] = salt_buf2[1];
w2_t[2] = salt_buf2[2];
w2_t[3] = salt_buf2[3];
w3_t[0] = salt_buf3[0];
w3_t[1] = salt_buf3[1];
w3_t[2] = salt_buf3[2];
w3_t[3] = salt_buf3[3];
u32x ipad[8]; u32x ipad[8];
u32x opad[8]; u32x opad[8];