arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-15 12:29:35 +00:00
Code Issues Releases Wiki Activity
153a8704e0
hashcat/OpenCL/m02710_a0-optimized.cl

783 lines
30 KiB
Common Lisp
Raw Normal View History

Initial commit
2015-12-04 14:47:52 +00:00
/**
Added docs/credits.txt Added docs/team.txt
2016-09-11 20:20:15 +00:00
* Author......: See docs/credits.txt
Initial commit
2015-12-04 14:47:52 +00:00
* License.....: MIT
*/
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
#define NEW_SIMD_CODE
Prepare OpenCL kernels for non-static compilation
2019-03-22 14:16:25 +00:00
#ifdef KERNEL_STATIC
Get rid of OpenCL/inc_hash_constants.h and OpenCL/inc_hash_functions.cl
2019-03-21 22:00:38 +00:00
#include "inc_vendor.h"
#include "inc_types.h"
Add missing inc_platform.h include
2019-04-26 11:59:43 +00:00
#include "inc_platform.cl"
Move files from include/ to OpenCL/ if they are used within kernels Rename includes in OpenCL so that it's easier to recognize them as such
2016-05-25 21:04:26 +00:00
#include "inc_common.cl"
Rename inc_rp.X to inc_rp_optimized.X
2017-08-11 09:25:47 +00:00
#include "inc_rp_optimized.h"
#include "inc_rp_optimized.cl"
Move files from include/ to OpenCL/ if they are used within kernels Rename includes in OpenCL so that it's easier to recognize them as such
2016-05-25 21:04:26 +00:00
#include "inc_simd.cl"
Get rid of OpenCL/inc_hash_constants.h and OpenCL/inc_hash_functions.cl
2019-03-21 22:00:38 +00:00
#include "inc_hash_md5.cl"
Prepare OpenCL kernels for non-static compilation
2019-03-22 14:16:25 +00:00
#endif
Initial commit
2015-12-04 14:47:52 +00:00
Prepare _a0 kernel for SIMD
2016-02-01 20:06:56 +00:00
#if VECT_SIZE == 1
Replace (u32x) (...) with make_u32x (...)
2019-05-08 13:21:22 +00:00
#define uint_to_hex_lower8(i) make_u32x (l_bin2asc[(i)])
Prepare _a0 kernel for SIMD
2016-02-01 20:06:56 +00:00
#elif VECT_SIZE == 2
Replace (u32x) (...) with make_u32x (...)
2019-05-08 13:21:22 +00:00
#define uint_to_hex_lower8(i) make_u32x (l_bin2asc[(i).s0], l_bin2asc[(i).s1])
Prepare _a0 kernel for SIMD
2016-02-01 20:06:56 +00:00
#elif VECT_SIZE == 4
Replace (u32x) (...) with make_u32x (...)
2019-05-08 13:21:22 +00:00
#define uint_to_hex_lower8(i) make_u32x (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3])
Prepare _a0 kernel for SIMD
2016-02-01 20:06:56 +00:00
#elif VECT_SIZE == 8
Replace (u32x) (...) with make_u32x (...)
2019-05-08 13:21:22 +00:00
#define uint_to_hex_lower8(i) make_u32x (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7])
Allow and support vector-width 16, which is current maximum for OpenCL. Closes #226.
2016-02-16 15:42:08 +00:00
#elif VECT_SIZE == 16
Replace (u32x) (...) with make_u32x (...)
2019-05-08 13:21:22 +00:00
#define uint_to_hex_lower8(i) make_u32x (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf])
Prepare _a0 kernel for SIMD
2016-02-01 20:06:56 +00:00
#endif
Initial commit
2015-12-04 14:47:52 +00:00
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
KERNEL_FQ void m02710_m04 (KERN_ATTR_RULES ())
Initial commit
2015-12-04 14:47:52 +00:00
{
/**
* modifier
*/
Switch the datatypes of the variables responsible for work-item count and work-item size from u32 to u64
2017-08-19 14:39:22 +00:00
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
Initial commit
2015-12-04 14:47:52 +00:00
/**
* bin2asc table
*/
Add LOCAL_VK to make use of __shared__
2019-05-07 07:01:32 +00:00
LOCAL_VK u32 l_bin2asc[256];
Initial commit
2015-12-04 14:47:52 +00:00
New Strategy: Instead of using volatile, mark the mode as unstable. Remove all volatiles
2019-03-03 18:18:56 +00:00
for (u32 i = lid; i < 256; i += lsz)
Upgrade kernel to support dynamic local work sizes
2016-01-19 15:06:03 +00:00
{
const u32 i0 = (i >> 0) & 15;
const u32 i1 = (i >> 4) & 15;
l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
}
Initial commit
2015-12-04 14:47:52 +00:00
Replace barrier() with SYNC_THREADS()
2019-04-26 11:34:07 +00:00
SYNC_THREADS ();
Initial commit
2015-12-04 14:47:52 +00:00
if (gid >= gid_max) return;
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
/**
* base
*/
u32 pw_buf0[4];
u32 pw_buf1[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];
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];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
/**
* salt
*/
u32 salt_buf0[4];
u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[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];
salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
salt_buf2[0] = 0;
salt_buf2[1] = 0;
salt_buf2[2] = 0;
salt_buf2[3] = 0;
salt_buf3[0] = 0;
salt_buf3[1] = 0;
salt_buf3[2] = 0;
salt_buf3[3] = 0;
const u32 salt_len = salt_bufs[salt_pos].salt_len;
Initial commit
2015-12-04 14:47:52 +00:00
/**
* loop
*/
Fixed two major problems 1) SIMD code for all attack-mode Macro vector_accessible() was not refactored and missing completely. Had to rename variables rules_cnt, combs_cnt and bfs_cnt into il_cnt which was a good thing anyway as with new SIMD code they all act in the same way. 2) SIMD code for attack-mode 0 With new SIMD code, apply_rules_vect() has to return u32 not u32x. This has massive impact on all *_a0 kernels. I've rewritten most of them. Deep testing using test.sh is still required. Some kernel need more fixes: - Some are kind of completely incompatible like m10400 but they still use old check_* includes, we should get rid of them as they are no longer neccessary as we have simd.c - Some have a chance but require additional effort like m11500. We can use commented out "#define NEW_SIMD_CODE" to find them This change can have negative impact on -a0 performance for device that require vectorization. That is mostly CPU devices. New GPU's are all scalar, so they wont get hurt by this. This change also proofes that there's no way to efficiently vectorize kernel rules with new SIMD code, but it enables the addition of the rule functions like @ that we were missing for some long time. This is a TODO.
2016-02-27 16:18:54 +00:00
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
Initial commit
2015-12-04 14:47:52 +00:00
{
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
u32x w0[4] = { 0 };
u32x w1[4] = { 0 };
u32x w2[4] = { 0 };
u32x w3[4] = { 0 };
Initial commit
2015-12-04 14:47:52 +00:00
Get rid of src/rp_kernel_on_cpu.c and src/rp_kernel_on_cpu_optimized.c and use OpenCL emulated kernel version
2019-03-25 11:24:04 +00:00
const u32x out_len = apply_rules_vect_optimized (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
Initial commit
2015-12-04 14:47:52 +00:00
Fixed two major problems 1) SIMD code for all attack-mode Macro vector_accessible() was not refactored and missing completely. Had to rename variables rules_cnt, combs_cnt and bfs_cnt into il_cnt which was a good thing anyway as with new SIMD code they all act in the same way. 2) SIMD code for attack-mode 0 With new SIMD code, apply_rules_vect() has to return u32 not u32x. This has massive impact on all *_a0 kernels. I've rewritten most of them. Deep testing using test.sh is still required. Some kernel need more fixes: - Some are kind of completely incompatible like m10400 but they still use old check_* includes, we should get rid of them as they are no longer neccessary as we have simd.c - Some have a chance but require additional effort like m11500. We can use commented out "#define NEW_SIMD_CODE" to find them This change can have negative impact on -a0 performance for device that require vectorization. That is mostly CPU devices. New GPU's are all scalar, so they wont get hurt by this. This change also proofes that there's no way to efficiently vectorize kernel rules with new SIMD code, but it enables the addition of the rule functions like @ that we were missing for some long time. This is a TODO.
2016-02-27 16:18:54 +00:00
append_0x80_2x4_VV (w0, w1, out_len);
Initial commit
2015-12-04 14:47:52 +00:00
w3[2] = out_len * 8;
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
w3[3] = 0;
Initial commit
2015-12-04 14:47:52 +00:00
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
u32x a = MD5M_A;
u32x b = MD5M_B;
u32x c = MD5M_C;
u32x d = MD5M_D;
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
2017-08-30 13:32:09 +00:00
u32x t;
MD5_STEP (MD5_H1, a, b, c, d, w1[1], MD5C20, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w2[0], MD5C21, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w2[3], MD5C22, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w3[2], MD5C23, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w0[1], MD5C24, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w1[0], MD5C25, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w1[3], MD5C26, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w2[2], MD5C27, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w3[1], MD5C28, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w0[0], MD5C29, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w0[3], MD5C2a, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w1[2], MD5C2b, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w2[1], MD5C2c, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w3[0], MD5C2d, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w3[3], MD5C2e, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w0[2], MD5C2f, MD5S23);
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
a += MD5M_A;
b += MD5M_B;
c += MD5M_C;
d += MD5M_D;
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
w0[0] = uint_to_hex_lower8 ((a >> 0) & 255) << 0
| uint_to_hex_lower8 ((a >> 8) & 255) << 16;
w0[1] = uint_to_hex_lower8 ((a >> 16) & 255) << 0
| uint_to_hex_lower8 ((a >> 24) & 255) << 16;
w0[2] = uint_to_hex_lower8 ((b >> 0) & 255) << 0
| uint_to_hex_lower8 ((b >> 8) & 255) << 16;
w0[3] = uint_to_hex_lower8 ((b >> 16) & 255) << 0
| uint_to_hex_lower8 ((b >> 24) & 255) << 16;
w1[0] = uint_to_hex_lower8 ((c >> 0) & 255) << 0
| uint_to_hex_lower8 ((c >> 8) & 255) << 16;
w1[1] = uint_to_hex_lower8 ((c >> 16) & 255) << 0
| uint_to_hex_lower8 ((c >> 24) & 255) << 16;
w1[2] = uint_to_hex_lower8 ((d >> 0) & 255) << 0
| uint_to_hex_lower8 ((d >> 8) & 255) << 16;
w1[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0
| uint_to_hex_lower8 ((d >> 24) & 255) << 16;
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
w2[0] = salt_buf0[0];
w2[1] = salt_buf0[1];
w2[2] = salt_buf0[2];
w2[3] = salt_buf0[3];
w3[0] = salt_buf1[0];
w3[1] = salt_buf1[1];
w3[2] = salt_buf1[2];
w3[3] = salt_buf1[3];
Initial commit
2015-12-04 14:47:52 +00:00
a = MD5M_A;
b = MD5M_B;
c = MD5M_C;
d = MD5M_D;
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
2017-08-30 13:32:09 +00:00
MD5_STEP (MD5_H1, a, b, c, d, w1[1], MD5C20, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w2[0], MD5C21, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w2[3], MD5C22, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w3[2], MD5C23, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w0[1], MD5C24, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w1[0], MD5C25, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w1[3], MD5C26, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w2[2], MD5C27, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w3[1], MD5C28, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w0[0], MD5C29, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w0[3], MD5C2a, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w1[2], MD5C2b, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w2[1], MD5C2c, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w3[0], MD5C2d, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w3[3], MD5C2e, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w0[2], MD5C2f, MD5S23);
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
const u32x r_a = a + MD5M_A;
const u32x r_b = b + MD5M_B;
const u32x r_c = c + MD5M_C;
const u32x r_d = d + MD5M_D;
Initial commit
2015-12-04 14:47:52 +00:00
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
const u32x r_14 = (32 + salt_len) * 8;
Initial commit
2015-12-04 14:47:52 +00:00
a = r_a;
b = r_b;
c = r_c;
d = r_d;
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C00, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C01, MD5S01);
MD5_STEP0(MD5_Fo, c, d, a, b, MD5C02, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C03, MD5S03);
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C04, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C05, MD5S01);
MD5_STEP0(MD5_Fo, c, d, a, b, MD5C06, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C07, MD5S03);
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C08, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C09, MD5S01);
MD5_STEP0(MD5_Fo, c, d, a, b, MD5C0a, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C0b, MD5S03);
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C0c, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, r_14, MD5C0e, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C0f, MD5S03);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C10, MD5S10);
MD5_STEP0(MD5_Go, d, a, b, c, MD5C11, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C12, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C13, MD5S13);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C14, MD5S10);
MD5_STEP0(MD5_Go, d, a, b, c, MD5C15, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C16, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C17, MD5S13);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, r_14, MD5C19, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C1a, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C1b, MD5S13);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C1c, MD5S10);
MD5_STEP0(MD5_Go, d, a, b, c, MD5C1d, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C1e, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C1f, MD5S13);
OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
2017-08-30 13:32:09 +00:00
MD5_STEP0(MD5_H1, a, b, c, d, MD5C20, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C21, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C22, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, r_14, MD5C23, MD5S23);
MD5_STEP0(MD5_H1, a, b, c, d, MD5C24, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C25, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C26, MD5S22);
MD5_STEP0(MD5_H2, b, c, d, a, MD5C27, MD5S23);
MD5_STEP0(MD5_H1, a, b, c, d, MD5C28, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C29, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C2a, MD5S22);
MD5_STEP0(MD5_H2, b, c, d, a, MD5C2b, MD5S23);
MD5_STEP0(MD5_H1, a, b, c, d, MD5C2c, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C2d, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C2e, MD5S22);
MD5_STEP0(MD5_H2, b, c, d, a, MD5C2f, MD5S23);
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP0(MD5_I , a, b, c, d, MD5C30, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, r_14, MD5C32, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C33, MD5S33);
MD5_STEP0(MD5_I , a, b, c, d, MD5C34, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C35, MD5S31);
MD5_STEP0(MD5_I , c, d, a, b, MD5C36, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C37, MD5S33);
MD5_STEP0(MD5_I , a, b, c, d, MD5C38, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C39, MD5S31);
MD5_STEP0(MD5_I , c, d, a, b, MD5C3a, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C3b, MD5S33);
MD5_STEP0(MD5_I , a, b, c, d, MD5C3c, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C3d, MD5S31);
MD5_STEP0(MD5_I , c, d, a, b, MD5C3e, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C3f, MD5S33);
a += r_a;
b += r_b;
c += r_c;
d += r_d;
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
COMPARE_M_SIMD (a, d, c, b);
Initial commit
2015-12-04 14:47:52 +00:00
}
}
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
KERNEL_FQ void m02710_m08 (KERN_ATTR_RULES ())
Initial commit
2015-12-04 14:47:52 +00:00
{
}
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
KERNEL_FQ void m02710_m16 (KERN_ATTR_RULES ())
Initial commit
2015-12-04 14:47:52 +00:00
{
}
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
KERNEL_FQ void m02710_s04 (KERN_ATTR_RULES ())
Initial commit
2015-12-04 14:47:52 +00:00
{
/**
* modifier
*/
Switch the datatypes of the variables responsible for work-item count and work-item size from u32 to u64
2017-08-19 14:39:22 +00:00
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
Initial commit
2015-12-04 14:47:52 +00:00
/**
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
* bin2asc table
Initial commit
2015-12-04 14:47:52 +00:00
*/
Add LOCAL_VK to make use of __shared__
2019-05-07 07:01:32 +00:00
LOCAL_VK u32 l_bin2asc[256];
Initial commit
2015-12-04 14:47:52 +00:00
New Strategy: Instead of using volatile, mark the mode as unstable. Remove all volatiles
2019-03-03 18:18:56 +00:00
for (u32 i = lid; i < 256; i += lsz)
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
{
const u32 i0 = (i >> 0) & 15;
const u32 i1 = (i >> 4) & 15;
l_bin2asc[i] = ((i0 < 10) ? '0' + i0 : 'a' - 10 + i0) << 8
| ((i1 < 10) ? '0' + i1 : 'a' - 10 + i1) << 0;
}
Replace barrier() with SYNC_THREADS()
2019-04-26 11:34:07 +00:00
SYNC_THREADS ();
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
if (gid >= gid_max) return;
/**
* base
*/
Initial commit
2015-12-04 14:47:52 +00:00
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
u32 pw_buf0[4];
- Dropped all vector code since new GPU's are all scalar, makes the code much easier - Some performance on low-end GPU may drop because of that, but only for a few hash-modes - Dropped scalar code (aka warp) since we do not have any vector datatypes anymore - Renamed C++ overloading functions memcat32_9 -> memcat_c32_w4x4_a3x4 - Still need to fix kernels to new function names, needs to be done manually - Temperature Management needs to be rewritten partially because of conflicting datatypes names - Added code to create different codepaths for NV on AMD in runtime in host (see data.vendor_id) - Added code to create different codepaths for NV on AMD in runtime in kernels (see IS_NV and IS_AMD) - First tests working for -m 0, for example - Great performance increases in general for NV so far - Tested amp_* and markov_* kernel - Migrated special NV optimizations for rule processor
2015-12-15 11:04:22 +00:00
u32 pw_buf1[4];
Initial commit
2015-12-04 14:47:52 +00:00
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
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];
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];
Initial commit
2015-12-04 14:47:52 +00:00
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Initial commit
2015-12-04 14:47:52 +00:00
/**
* salt
*/
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
u32 salt_buf0[4];
u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[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];
salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
salt_buf2[0] = 0;
salt_buf2[1] = 0;
salt_buf2[2] = 0;
salt_buf2[3] = 0;
salt_buf3[0] = 0;
salt_buf3[1] = 0;
salt_buf3[2] = 0;
salt_buf3[3] = 0;
Initial commit
2015-12-04 14:47:52 +00:00
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
const u32 salt_len = salt_bufs[salt_pos].salt_len;
Initial commit
2015-12-04 14:47:52 +00:00
/**
* digest
*/
const u32 search[4] =
{
digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2],
digests_buf[digests_offset].digest_buf[DGST_R3]
};
/**
* loop
*/
Fixed two major problems 1) SIMD code for all attack-mode Macro vector_accessible() was not refactored and missing completely. Had to rename variables rules_cnt, combs_cnt and bfs_cnt into il_cnt which was a good thing anyway as with new SIMD code they all act in the same way. 2) SIMD code for attack-mode 0 With new SIMD code, apply_rules_vect() has to return u32 not u32x. This has massive impact on all *_a0 kernels. I've rewritten most of them. Deep testing using test.sh is still required. Some kernel need more fixes: - Some are kind of completely incompatible like m10400 but they still use old check_* includes, we should get rid of them as they are no longer neccessary as we have simd.c - Some have a chance but require additional effort like m11500. We can use commented out "#define NEW_SIMD_CODE" to find them This change can have negative impact on -a0 performance for device that require vectorization. That is mostly CPU devices. New GPU's are all scalar, so they wont get hurt by this. This change also proofes that there's no way to efficiently vectorize kernel rules with new SIMD code, but it enables the addition of the rule functions like @ that we were missing for some long time. This is a TODO.
2016-02-27 16:18:54 +00:00
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
Initial commit
2015-12-04 14:47:52 +00:00
{
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
u32x w0[4] = { 0 };
u32x w1[4] = { 0 };
u32x w2[4] = { 0 };
u32x w3[4] = { 0 };
Initial commit
2015-12-04 14:47:52 +00:00
Get rid of src/rp_kernel_on_cpu.c and src/rp_kernel_on_cpu_optimized.c and use OpenCL emulated kernel version
2019-03-25 11:24:04 +00:00
const u32x out_len = apply_rules_vect_optimized (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
Initial commit
2015-12-04 14:47:52 +00:00
Fixed two major problems 1) SIMD code for all attack-mode Macro vector_accessible() was not refactored and missing completely. Had to rename variables rules_cnt, combs_cnt and bfs_cnt into il_cnt which was a good thing anyway as with new SIMD code they all act in the same way. 2) SIMD code for attack-mode 0 With new SIMD code, apply_rules_vect() has to return u32 not u32x. This has massive impact on all *_a0 kernels. I've rewritten most of them. Deep testing using test.sh is still required. Some kernel need more fixes: - Some are kind of completely incompatible like m10400 but they still use old check_* includes, we should get rid of them as they are no longer neccessary as we have simd.c - Some have a chance but require additional effort like m11500. We can use commented out "#define NEW_SIMD_CODE" to find them This change can have negative impact on -a0 performance for device that require vectorization. That is mostly CPU devices. New GPU's are all scalar, so they wont get hurt by this. This change also proofes that there's no way to efficiently vectorize kernel rules with new SIMD code, but it enables the addition of the rule functions like @ that we were missing for some long time. This is a TODO.
2016-02-27 16:18:54 +00:00
append_0x80_2x4_VV (w0, w1, out_len);
Initial commit
2015-12-04 14:47:52 +00:00
w3[2] = out_len * 8;
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
w3[3] = 0;
Initial commit
2015-12-04 14:47:52 +00:00
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
u32x a = MD5M_A;
u32x b = MD5M_B;
u32x c = MD5M_C;
u32x d = MD5M_D;
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
2017-08-30 13:32:09 +00:00
u32x t;
MD5_STEP (MD5_H1, a, b, c, d, w1[1], MD5C20, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w2[0], MD5C21, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w2[3], MD5C22, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w3[2], MD5C23, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w0[1], MD5C24, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w1[0], MD5C25, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w1[3], MD5C26, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w2[2], MD5C27, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w3[1], MD5C28, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w0[0], MD5C29, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w0[3], MD5C2a, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w1[2], MD5C2b, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w2[1], MD5C2c, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w3[0], MD5C2d, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w3[3], MD5C2e, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w0[2], MD5C2f, MD5S23);
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
a += MD5M_A;
b += MD5M_B;
c += MD5M_C;
d += MD5M_D;
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
w0[0] = uint_to_hex_lower8 ((a >> 0) & 255) << 0
| uint_to_hex_lower8 ((a >> 8) & 255) << 16;
w0[1] = uint_to_hex_lower8 ((a >> 16) & 255) << 0
| uint_to_hex_lower8 ((a >> 24) & 255) << 16;
w0[2] = uint_to_hex_lower8 ((b >> 0) & 255) << 0
| uint_to_hex_lower8 ((b >> 8) & 255) << 16;
w0[3] = uint_to_hex_lower8 ((b >> 16) & 255) << 0
| uint_to_hex_lower8 ((b >> 24) & 255) << 16;
w1[0] = uint_to_hex_lower8 ((c >> 0) & 255) << 0
| uint_to_hex_lower8 ((c >> 8) & 255) << 16;
w1[1] = uint_to_hex_lower8 ((c >> 16) & 255) << 0
| uint_to_hex_lower8 ((c >> 24) & 255) << 16;
w1[2] = uint_to_hex_lower8 ((d >> 0) & 255) << 0
| uint_to_hex_lower8 ((d >> 8) & 255) << 16;
w1[3] = uint_to_hex_lower8 ((d >> 16) & 255) << 0
| uint_to_hex_lower8 ((d >> 24) & 255) << 16;
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
w2[0] = salt_buf0[0];
w2[1] = salt_buf0[1];
w2[2] = salt_buf0[2];
w2[3] = salt_buf0[3];
w3[0] = salt_buf1[0];
w3[1] = salt_buf1[1];
w3[2] = salt_buf1[2];
w3[3] = salt_buf1[3];
Initial commit
2015-12-04 14:47:52 +00:00
a = MD5M_A;
b = MD5M_B;
c = MD5M_C;
d = MD5M_D;
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
2017-08-30 13:32:09 +00:00
MD5_STEP (MD5_H1, a, b, c, d, w1[1], MD5C20, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w2[0], MD5C21, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w2[3], MD5C22, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w3[2], MD5C23, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w0[1], MD5C24, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w1[0], MD5C25, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w1[3], MD5C26, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w2[2], MD5C27, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w3[1], MD5C28, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w0[0], MD5C29, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w0[3], MD5C2a, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w1[2], MD5C2b, MD5S23);
MD5_STEP (MD5_H1, a, b, c, d, w2[1], MD5C2c, MD5S20);
MD5_STEP (MD5_H2, d, a, b, c, w3[0], MD5C2d, MD5S21);
MD5_STEP (MD5_H1, c, d, a, b, w3[3], MD5C2e, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, w0[2], MD5C2f, MD5S23);
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
const u32x r_a = a + MD5M_A;
const u32x r_b = b + MD5M_B;
const u32x r_c = c + MD5M_C;
const u32x r_d = d + MD5M_D;
Initial commit
2015-12-04 14:47:52 +00:00
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
const u32x r_14 = (32 + salt_len) * 8;
Initial commit
2015-12-04 14:47:52 +00:00
a = r_a;
b = r_b;
c = r_c;
d = r_d;
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C00, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C01, MD5S01);
MD5_STEP0(MD5_Fo, c, d, a, b, MD5C02, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C03, MD5S03);
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C04, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C05, MD5S01);
MD5_STEP0(MD5_Fo, c, d, a, b, MD5C06, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C07, MD5S03);
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C08, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C09, MD5S01);
MD5_STEP0(MD5_Fo, c, d, a, b, MD5C0a, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C0b, MD5S03);
MD5_STEP0(MD5_Fo, a, b, c, d, MD5C0c, MD5S00);
MD5_STEP0(MD5_Fo, d, a, b, c, MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, r_14, MD5C0e, MD5S02);
MD5_STEP0(MD5_Fo, b, c, d, a, MD5C0f, MD5S03);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C10, MD5S10);
MD5_STEP0(MD5_Go, d, a, b, c, MD5C11, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C12, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C13, MD5S13);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C14, MD5S10);
MD5_STEP0(MD5_Go, d, a, b, c, MD5C15, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C16, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C17, MD5S13);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, r_14, MD5C19, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C1a, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C1b, MD5S13);
MD5_STEP0(MD5_Go, a, b, c, d, MD5C1c, MD5S10);
MD5_STEP0(MD5_Go, d, a, b, c, MD5C1d, MD5S11);
MD5_STEP0(MD5_Go, c, d, a, b, MD5C1e, MD5S12);
MD5_STEP0(MD5_Go, b, c, d, a, MD5C1f, MD5S13);
OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
2017-08-30 13:32:09 +00:00
MD5_STEP0(MD5_H1, a, b, c, d, MD5C20, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C21, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C22, MD5S22);
MD5_STEP (MD5_H2, b, c, d, a, r_14, MD5C23, MD5S23);
MD5_STEP0(MD5_H1, a, b, c, d, MD5C24, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C25, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C26, MD5S22);
MD5_STEP0(MD5_H2, b, c, d, a, MD5C27, MD5S23);
MD5_STEP0(MD5_H1, a, b, c, d, MD5C28, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C29, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C2a, MD5S22);
MD5_STEP0(MD5_H2, b, c, d, a, MD5C2b, MD5S23);
MD5_STEP0(MD5_H1, a, b, c, d, MD5C2c, MD5S20);
MD5_STEP0(MD5_H2, d, a, b, c, MD5C2d, MD5S21);
MD5_STEP0(MD5_H1, c, d, a, b, MD5C2e, MD5S22);
MD5_STEP0(MD5_H2, b, c, d, a, MD5C2f, MD5S23);
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP0(MD5_I , a, b, c, d, MD5C30, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, r_14, MD5C32, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C33, MD5S33);
MD5_STEP0(MD5_I , a, b, c, d, MD5C34, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C35, MD5S31);
MD5_STEP0(MD5_I , c, d, a, b, MD5C36, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C37, MD5S33);
MD5_STEP0(MD5_I , a, b, c, d, MD5C38, MD5S30);
MD5_STEP0(MD5_I , d, a, b, c, MD5C39, MD5S31);
MD5_STEP0(MD5_I , c, d, a, b, MD5C3a, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C3b, MD5S33);
MD5_STEP0(MD5_I , a, b, c, d, MD5C3c, MD5S30);
Cleanup -m 27xx kernels to latest standard
2016-03-19 16:07:28 +00:00
if (MATCHES_NONE_VS ((a + r_a), search[0])) continue;
Initial commit
2015-12-04 14:47:52 +00:00
MD5_STEP0(MD5_I , d, a, b, c, MD5C3d, MD5S31);
MD5_STEP0(MD5_I , c, d, a, b, MD5C3e, MD5S32);
MD5_STEP0(MD5_I , b, c, d, a, MD5C3f, MD5S33);
a += r_a;
b += r_b;
c += r_c;
d += r_d;
Converted to new SIMD: -m 2710 -a 0
2016-02-05 20:20:42 +00:00
COMPARE_S_SIMD (a, d, c, b);
Initial commit
2015-12-04 14:47:52 +00:00
}
}
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
KERNEL_FQ void m02710_s08 (KERN_ATTR_RULES ())
Initial commit
2015-12-04 14:47:52 +00:00
{
}
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 21:27:58 +00:00
KERNEL_FQ void m02710_s16 (KERN_ATTR_RULES ())
Initial commit
2015-12-04 14:47:52 +00:00
{
}
Reference in New Issue Copy Permalink