From 02e2279d594c57cb857201a825e06624f0ad3a5b Mon Sep 17 00:00:00 2001
From: jsteube <jens.steube@gmail.com>
Date: Mon, 24 Jul 2017 14:33:34 +0200
Subject: [PATCH] Optimized -m 8500 for ROCm
---
OpenCL/m08500_a0-optimized.cl | 36 +++++++++---------------------
OpenCL/m08500_a1-optimized.cl | 38 +++++++++----------------------
OpenCL/m08500_a3-optimized.cl | 42 +++++++++++++----------------------
3 files changed, 37 insertions(+), 79 deletions(-)
diff --git a/OpenCL/m08500_a0-optimized.cl b/OpenCL/m08500_a0-optimized.cl
index ba9bc66fd..5626b51ce 100644
--- a/OpenCL/m08500_a0-optimized.cl
+++ b/OpenCL/m08500_a0-optimized.cl
@@ -508,17 +508,17 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
}
}
-void transform_racf_key (const u32x w0, const u32x w1, u32x key[2], __local u32 *s_ascii_to_ebcdic_pc)
+void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
{
- key[0] = BOX1 (((w0 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0
- | BOX1 (((w0 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8
- | BOX1 (((w0 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16
- | BOX1 (((w0 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24;
+ key[0] = BOX1 (((w0 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
+ | BOX1 (((w0 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
+ | BOX1 (((w0 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
+ | BOX1 (((w0 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
- key[1] = BOX1 (((w1 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0
- | BOX1 (((w1 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8
- | BOX1 (((w1 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16
- | BOX1 (((w1 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24;
+ key[1] = BOX1 (((w1 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
+ | BOX1 (((w1 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
+ | BOX1 (((w1 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
+ | BOX1 (((w1 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
}
__kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
@@ -535,7 +535,6 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
- __local u32 s_ascii_to_ebcdic_pc[256];
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
@@ -560,13 +559,6 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
s_skb[7][i] = c_skb[7][i];
}
- for (u32 i = lid; i < 256; i += lsz)
- {
- s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i];
- }
-
- barrier (CLK_LOCAL_MEM_FENCE);
-
if (gid >= gid_max) return;
/**
@@ -615,7 +607,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x key[2];
- transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc);
+ transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];
@@ -662,7 +654,6 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
- __local u32 s_ascii_to_ebcdic_pc[256];
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
@@ -687,11 +678,6 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
s_skb[7][i] = c_skb[7][i];
}
- for (u32 i = lid; i < 256; i += lsz)
- {
- s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i];
- }
-
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
@@ -754,7 +740,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x key[2];
- transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc);
+ transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];
diff --git a/OpenCL/m08500_a1-optimized.cl b/OpenCL/m08500_a1-optimized.cl
index f76728dd8..6b7dca17e 100644
--- a/OpenCL/m08500_a1-optimized.cl
+++ b/OpenCL/m08500_a1-optimized.cl
@@ -506,17 +506,17 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
}
}
-void transform_racf_key (const u32x w0, const u32x w1, u32x key[2], __local u32 *s_ascii_to_ebcdic_pc)
+void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
{
- key[0] = BOX1 (((w0 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0
- | BOX1 (((w0 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8
- | BOX1 (((w0 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16
- | BOX1 (((w0 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24;
+ key[0] = BOX1 (((w0 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
+ | BOX1 (((w0 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
+ | BOX1 (((w0 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
+ | BOX1 (((w0 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
- key[1] = BOX1 (((w1 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0
- | BOX1 (((w1 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8
- | BOX1 (((w1 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16
- | BOX1 (((w1 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24;
+ key[1] = BOX1 (((w1 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
+ | BOX1 (((w1 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
+ | BOX1 (((w1 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
+ | BOX1 (((w1 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
}
__kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
@@ -533,7 +533,6 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
- __local u32 s_ascii_to_ebcdic_pc[256];
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
@@ -558,13 +557,6 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
s_skb[7][i] = c_skb[7][i];
}
- for (u32 i = lid; i < 256; i += lsz)
- {
- s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i];
- }
-
- barrier (CLK_LOCAL_MEM_FENCE);
-
if (gid >= gid_max) return;
/**
@@ -656,7 +648,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x key[2];
- transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc);
+ transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];
@@ -703,7 +695,6 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
- __local u32 s_ascii_to_ebcdic_pc[256];
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
@@ -728,13 +719,6 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
s_skb[7][i] = c_skb[7][i];
}
- for (u32 i = lid; i < 256; i += lsz)
- {
- s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i];
- }
-
- barrier (CLK_LOCAL_MEM_FENCE);
-
if (gid >= gid_max) return;
/**
@@ -838,7 +822,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x key[2];
- transform_racf_key (w0[0], w0[1], key, s_ascii_to_ebcdic_pc);
+ transform_racf_key (w0[0], w0[1], key);
const u32x c = key[0];
const u32x d = key[1];
diff --git a/OpenCL/m08500_a3-optimized.cl b/OpenCL/m08500_a3-optimized.cl
index 3bdeae561..8cff7bcc2 100644
--- a/OpenCL/m08500_a3-optimized.cl
+++ b/OpenCL/m08500_a3-optimized.cl
@@ -506,20 +506,20 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
}
}
-void transform_racf_key (const u32x w0, const u32x w1, u32x key[2], __local u32 *s_ascii_to_ebcdic_pc)
+void transform_racf_key (const u32x w0, const u32x w1, u32x key[2])
{
- key[0] = BOX1 (((w0 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0
- | BOX1 (((w0 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8
- | BOX1 (((w0 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16
- | BOX1 (((w0 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24;
+ key[0] = BOX1 (((w0 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
+ | BOX1 (((w0 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
+ | BOX1 (((w0 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
+ | BOX1 (((w0 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
- key[1] = BOX1 (((w1 >> 0) & 0xff), s_ascii_to_ebcdic_pc) << 0
- | BOX1 (((w1 >> 8) & 0xff), s_ascii_to_ebcdic_pc) << 8
- | BOX1 (((w1 >> 16) & 0xff), s_ascii_to_ebcdic_pc) << 16
- | BOX1 (((w1 >> 24) & 0xff), s_ascii_to_ebcdic_pc) << 24;
+ key[1] = BOX1 (((w1 >> 0) & 0xff), c_ascii_to_ebcdic_pc) << 0
+ | BOX1 (((w1 >> 8) & 0xff), c_ascii_to_ebcdic_pc) << 8
+ | BOX1 (((w1 >> 16) & 0xff), c_ascii_to_ebcdic_pc) << 16
+ | BOX1 (((w1 >> 24) & 0xff), c_ascii_to_ebcdic_pc) << 24;
}
-void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u32 *s_ascii_to_ebcdic_pc, u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
+void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{
/**
* modifier
@@ -557,7 +557,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u3
u32x key[2];
- transform_racf_key (w0, w1, key, s_ascii_to_ebcdic_pc);
+ transform_racf_key (w0, w1, key);
const u32x c = key[0];
const u32x d = key[1];
@@ -582,7 +582,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u3
}
}
-void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u32 *s_ascii_to_ebcdic_pc, u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
+void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset)
{
/**
* modifier
@@ -632,7 +632,7 @@ void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], __local u3
u32x key[2];
- transform_racf_key (w0, w1, key, s_ascii_to_ebcdic_pc);
+ transform_racf_key (w0, w1, key);
const u32x c = key[0];
const u32x d = key[1];
@@ -671,7 +671,6 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
- __local u32 s_ascii_to_ebcdic_pc[256];
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
@@ -696,11 +695,6 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
s_skb[7][i] = c_skb[7][i];
}
- for (u32 i = lid; i < 256; i += lsz)
- {
- s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i];
- }
-
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
@@ -734,7 +728,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
- m08500m (s_SPtrans, s_skb, s_ascii_to_ebcdic_pc, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
+ m08500m (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
__kernel void m08500_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
@@ -759,7 +753,6 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* shared
*/
- __local u32 s_ascii_to_ebcdic_pc[256];
__local u32 s_SPtrans[8][64];
__local u32 s_skb[8][64];
@@ -784,11 +777,6 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
s_skb[7][i] = c_skb[7][i];
}
- for (u32 i = lid; i < 256; i += lsz)
- {
- s_ascii_to_ebcdic_pc[i] = c_ascii_to_ebcdic_pc[i];
- }
-
barrier (CLK_LOCAL_MEM_FENCE);
if (gid >= gid_max) return;
@@ -822,7 +810,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* main
*/
- m08500s (s_SPtrans, s_skb, s_ascii_to_ebcdic_pc, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
+ m08500s (s_SPtrans, s_skb, w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset);
}
__kernel void m08500_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)