From 0783289e2fd3015ff9be2d893df76d7950482b54 Mon Sep 17 00:00:00 2001 From: jsteube Date: Wed, 23 Aug 2017 13:40:22 +0200 Subject: [PATCH] Optimized a0 pure kernel for AMD --- OpenCL/inc_rp.h | 15 +++++++++++++++ OpenCL/m00000_a0.cl | 8 ++++---- OpenCL/m00010_a0.cl | 8 ++++---- OpenCL/m00020_a0.cl | 8 ++++---- OpenCL/m00030_a0.cl | 8 ++++---- OpenCL/m00040_a0.cl | 8 ++++---- OpenCL/m00050_a0.cl | 8 ++++---- OpenCL/m00060_a0.cl | 8 ++++---- OpenCL/m00100_a0.cl | 8 ++++---- OpenCL/m00110_a0.cl | 8 ++++---- OpenCL/m00120_a0.cl | 8 ++++---- OpenCL/m00130_a0.cl | 8 ++++---- OpenCL/m00140_a0.cl | 8 ++++---- OpenCL/m00150_a0.cl | 8 ++++---- OpenCL/m00160_a0.cl | 8 ++++---- OpenCL/m00300_a0.cl | 8 ++++---- OpenCL/m00900_a0.cl | 8 ++++---- OpenCL/m01000_a0.cl | 8 ++++---- OpenCL/m01100_a0.cl | 8 ++++---- OpenCL/m01300_a0.cl | 8 ++++---- OpenCL/m01400_a0.cl | 8 ++++---- OpenCL/m01410_a0.cl | 8 ++++---- OpenCL/m01420_a0.cl | 8 ++++---- OpenCL/m01430_a0.cl | 8 ++++---- OpenCL/m01440_a0.cl | 8 ++++---- OpenCL/m01450_a0.cl | 8 ++++---- OpenCL/m01460_a0.cl | 8 ++++---- OpenCL/m01700_a0.cl | 8 ++++---- OpenCL/m01710_a0.cl | 8 ++++---- OpenCL/m01720_a0.cl | 8 ++++---- OpenCL/m01730_a0.cl | 8 ++++---- OpenCL/m01740_a0.cl | 8 ++++---- OpenCL/m01750_a0.cl | 8 ++++---- OpenCL/m01760_a0.cl | 8 ++++---- OpenCL/m02610_a0.cl | 8 ++++---- OpenCL/m02810_a0.cl | 8 ++++---- OpenCL/m03710_a0.cl | 8 ++++---- OpenCL/m03800_a0.cl | 8 ++++---- OpenCL/m03910_a0.cl | 8 ++++---- OpenCL/m04010_a0.cl | 8 ++++---- OpenCL/m04110_a0.cl | 8 ++++---- OpenCL/m04310_a0.cl | 8 ++++---- OpenCL/m04400_a0.cl | 8 ++++---- OpenCL/m04500_a0.cl | 8 ++++---- OpenCL/m04520_a0.cl | 8 ++++---- OpenCL/m04700_a0.cl | 8 ++++---- OpenCL/m04800_a0.cl | 8 ++++---- OpenCL/m04900_a0.cl | 8 ++++---- OpenCL/m05100_a0.cl | 8 ++++---- OpenCL/m05300_a0.cl | 8 ++++---- OpenCL/m05400_a0.cl | 8 ++++---- OpenCL/m05500_a0.cl | 8 ++++---- OpenCL/m05600_a0.cl | 8 ++++---- OpenCL/m06000_a0.cl | 8 ++++---- OpenCL/m06100_a0.cl | 4 ++-- OpenCL/m07000_a0.cl | 8 ++++---- OpenCL/m07300_a0.cl | 8 ++++---- OpenCL/m07500_a0.cl | 8 ++++---- OpenCL/m08100_a0.cl | 8 ++++---- OpenCL/m08300_a0.cl | 8 ++++---- OpenCL/m08400_a0.cl | 8 ++++---- OpenCL/m09900_a0.cl | 8 ++++---- OpenCL/m10800_a0.cl | 4 ++-- OpenCL/m11000_a0.cl | 8 ++++---- OpenCL/m11100_a0.cl | 8 ++++---- OpenCL/m11200_a0.cl | 8 ++++---- OpenCL/m11400_a0.cl | 8 ++++---- OpenCL/m12600_a0.cl | 8 ++++---- OpenCL/m13100_a0.cl | 8 ++++---- OpenCL/m13300_a0.cl | 8 ++++---- OpenCL/m13500_a0.cl | 8 ++++---- OpenCL/m13800_a0.cl | 8 ++++---- OpenCL/m13900_a0.cl | 8 ++++---- OpenCL/m14400_a0.cl | 8 ++++---- OpenCL/m15000_a0.cl | 8 ++++---- OpenCL/m15500_a0.cl | 8 ++++---- 76 files changed, 311 insertions(+), 296 deletions(-) diff --git a/OpenCL/inc_rp.h b/OpenCL/inc_rp.h index 4694fb19e..25fd41cda 100644 --- a/OpenCL/inc_rp.h +++ b/OpenCL/inc_rp.h @@ -58,3 +58,18 @@ #define RULE_OP_MANGLE_TITLE 'E' #define RP_PASSWORD_SIZE 256 + +#ifdef IS_NV +#define COPY_PW(x) \ + pw_t pw = (x); +#else +#define COPY_PW(x) \ + __local pw_t s_pws[64]; \ + s_pws[get_local_id(0)] = (x); +#endif + +#ifdef IS_NV +#define PASTE_PW pw; +#else +#define PASTE_PW s_pws[get_local_id(0)]; +#endif diff --git a/OpenCL/m00000_a0.cl b/OpenCL/m00000_a0.cl index 38a8404e0..cf4c711ed 100644 --- a/OpenCL/m00000_a0.cl +++ b/OpenCL/m00000_a0.cl @@ -30,7 +30,7 @@ __kernel void m00000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m00000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m00000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m00000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00010_a0.cl b/OpenCL/m00010_a0.cl index 307eaa215..2b9c28328 100644 --- a/OpenCL/m00010_a0.cl +++ b/OpenCL/m00010_a0.cl @@ -30,7 +30,7 @@ __kernel void m00010_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m00010_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m00010_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m00010_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00020_a0.cl b/OpenCL/m00020_a0.cl index bf38bf75b..a4e690cda 100644 --- a/OpenCL/m00020_a0.cl +++ b/OpenCL/m00020_a0.cl @@ -30,7 +30,7 @@ __kernel void m00020_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m00020_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m00020_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m00020_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00030_a0.cl b/OpenCL/m00030_a0.cl index 5361254d0..1f63bcca4 100644 --- a/OpenCL/m00030_a0.cl +++ b/OpenCL/m00030_a0.cl @@ -30,7 +30,7 @@ __kernel void m00030_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m00030_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m00030_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m00030_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00040_a0.cl b/OpenCL/m00040_a0.cl index 527d3bf86..213acb92d 100644 --- a/OpenCL/m00040_a0.cl +++ b/OpenCL/m00040_a0.cl @@ -30,7 +30,7 @@ __kernel void m00040_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m00040_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m00040_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m00040_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00050_a0.cl b/OpenCL/m00050_a0.cl index fb213928a..fa1a00ef9 100644 --- a/OpenCL/m00050_a0.cl +++ b/OpenCL/m00050_a0.cl @@ -30,7 +30,7 @@ __kernel void m00050_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m00050_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -95,7 +95,7 @@ __kernel void m00050_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -112,7 +112,7 @@ __kernel void m00050_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00060_a0.cl b/OpenCL/m00060_a0.cl index 9141e150f..e2047cd62 100644 --- a/OpenCL/m00060_a0.cl +++ b/OpenCL/m00060_a0.cl @@ -30,7 +30,7 @@ __kernel void m00060_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -51,7 +51,7 @@ __kernel void m00060_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m00060_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -118,7 +118,7 @@ __kernel void m00060_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00100_a0.cl b/OpenCL/m00100_a0.cl index 5cb60b354..c6979b3e7 100644 --- a/OpenCL/m00100_a0.cl +++ b/OpenCL/m00100_a0.cl @@ -30,7 +30,7 @@ __kernel void m00100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m00100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m00100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m00100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00110_a0.cl b/OpenCL/m00110_a0.cl index 1071fa7f2..ee869a8ed 100644 --- a/OpenCL/m00110_a0.cl +++ b/OpenCL/m00110_a0.cl @@ -30,7 +30,7 @@ __kernel void m00110_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m00110_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m00110_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m00110_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00120_a0.cl b/OpenCL/m00120_a0.cl index 0241c3a99..033a8302f 100644 --- a/OpenCL/m00120_a0.cl +++ b/OpenCL/m00120_a0.cl @@ -30,7 +30,7 @@ __kernel void m00120_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m00120_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m00120_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m00120_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00130_a0.cl b/OpenCL/m00130_a0.cl index 701792042..1195d5e2d 100644 --- a/OpenCL/m00130_a0.cl +++ b/OpenCL/m00130_a0.cl @@ -30,7 +30,7 @@ __kernel void m00130_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m00130_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m00130_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m00130_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00140_a0.cl b/OpenCL/m00140_a0.cl index 9d6ec9c2a..cf7e1627d 100644 --- a/OpenCL/m00140_a0.cl +++ b/OpenCL/m00140_a0.cl @@ -30,7 +30,7 @@ __kernel void m00140_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m00140_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m00140_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m00140_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00150_a0.cl b/OpenCL/m00150_a0.cl index ef0e756cf..0f4f9a060 100644 --- a/OpenCL/m00150_a0.cl +++ b/OpenCL/m00150_a0.cl @@ -30,7 +30,7 @@ __kernel void m00150_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m00150_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -95,7 +95,7 @@ __kernel void m00150_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -112,7 +112,7 @@ __kernel void m00150_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00160_a0.cl b/OpenCL/m00160_a0.cl index b5d19e28e..2e8422ac3 100644 --- a/OpenCL/m00160_a0.cl +++ b/OpenCL/m00160_a0.cl @@ -30,7 +30,7 @@ __kernel void m00160_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -51,7 +51,7 @@ __kernel void m00160_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m00160_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -118,7 +118,7 @@ __kernel void m00160_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00300_a0.cl b/OpenCL/m00300_a0.cl index 146273845..d1dfb8557 100644 --- a/OpenCL/m00300_a0.cl +++ b/OpenCL/m00300_a0.cl @@ -30,7 +30,7 @@ __kernel void m00300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m00300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -109,7 +109,7 @@ __kernel void m00300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -117,7 +117,7 @@ __kernel void m00300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m00900_a0.cl b/OpenCL/m00900_a0.cl index d595da7fb..516e73923 100644 --- a/OpenCL/m00900_a0.cl +++ b/OpenCL/m00900_a0.cl @@ -30,7 +30,7 @@ __kernel void m00900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m00900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m00900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m00900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01000_a0.cl b/OpenCL/m01000_a0.cl index 6019ef934..2fe1a2969 100644 --- a/OpenCL/m01000_a0.cl +++ b/OpenCL/m01000_a0.cl @@ -30,7 +30,7 @@ __kernel void m01000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m01000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m01000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m01000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01100_a0.cl b/OpenCL/m01100_a0.cl index 1a511754b..d79dc18da 100644 --- a/OpenCL/m01100_a0.cl +++ b/OpenCL/m01100_a0.cl @@ -30,7 +30,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -110,7 +110,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -127,7 +127,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01300_a0.cl b/OpenCL/m01300_a0.cl index 5eb7ad5c8..a4a4ec299 100644 --- a/OpenCL/m01300_a0.cl +++ b/OpenCL/m01300_a0.cl @@ -30,7 +30,7 @@ __kernel void m01300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m01300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m01300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m01300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01400_a0.cl b/OpenCL/m01400_a0.cl index 8ede23992..d7df4e289 100644 --- a/OpenCL/m01400_a0.cl +++ b/OpenCL/m01400_a0.cl @@ -30,7 +30,7 @@ __kernel void m01400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m01400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m01400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m01400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01410_a0.cl b/OpenCL/m01410_a0.cl index 75701e85f..d21a5195f 100644 --- a/OpenCL/m01410_a0.cl +++ b/OpenCL/m01410_a0.cl @@ -30,7 +30,7 @@ __kernel void m01410_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01410_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m01410_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m01410_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01420_a0.cl b/OpenCL/m01420_a0.cl index 6dae5250d..176a69ebf 100644 --- a/OpenCL/m01420_a0.cl +++ b/OpenCL/m01420_a0.cl @@ -30,7 +30,7 @@ __kernel void m01420_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha256_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m01420_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m01420_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha256_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m01420_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01430_a0.cl b/OpenCL/m01430_a0.cl index 8fbc16823..46e69addb 100644 --- a/OpenCL/m01430_a0.cl +++ b/OpenCL/m01430_a0.cl @@ -30,7 +30,7 @@ __kernel void m01430_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01430_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m01430_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m01430_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01440_a0.cl b/OpenCL/m01440_a0.cl index 064967b44..9bc6b20f7 100644 --- a/OpenCL/m01440_a0.cl +++ b/OpenCL/m01440_a0.cl @@ -30,7 +30,7 @@ __kernel void m01440_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha256_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m01440_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m01440_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha256_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m01440_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01450_a0.cl b/OpenCL/m01450_a0.cl index 93d5a5d0b..b54ad3748 100644 --- a/OpenCL/m01450_a0.cl +++ b/OpenCL/m01450_a0.cl @@ -30,7 +30,7 @@ __kernel void m01450_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01450_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -95,7 +95,7 @@ __kernel void m01450_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -112,7 +112,7 @@ __kernel void m01450_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01460_a0.cl b/OpenCL/m01460_a0.cl index a534a15c7..eb8556af9 100644 --- a/OpenCL/m01460_a0.cl +++ b/OpenCL/m01460_a0.cl @@ -30,7 +30,7 @@ __kernel void m01460_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -51,7 +51,7 @@ __kernel void m01460_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m01460_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -118,7 +118,7 @@ __kernel void m01460_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01700_a0.cl b/OpenCL/m01700_a0.cl index e0870abd2..65f11f602 100644 --- a/OpenCL/m01700_a0.cl +++ b/OpenCL/m01700_a0.cl @@ -30,7 +30,7 @@ __kernel void m01700_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m01700_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m01700_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m01700_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01710_a0.cl b/OpenCL/m01710_a0.cl index e5cda2fbb..8f74d7c47 100644 --- a/OpenCL/m01710_a0.cl +++ b/OpenCL/m01710_a0.cl @@ -30,7 +30,7 @@ __kernel void m01710_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01710_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m01710_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m01710_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01720_a0.cl b/OpenCL/m01720_a0.cl index 2b7fb4380..f2663f3d7 100644 --- a/OpenCL/m01720_a0.cl +++ b/OpenCL/m01720_a0.cl @@ -30,7 +30,7 @@ __kernel void m01720_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha512_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m01720_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m01720_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha512_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m01720_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01730_a0.cl b/OpenCL/m01730_a0.cl index 623ba72df..30d2a00a9 100644 --- a/OpenCL/m01730_a0.cl +++ b/OpenCL/m01730_a0.cl @@ -30,7 +30,7 @@ __kernel void m01730_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01730_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m01730_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m01730_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01740_a0.cl b/OpenCL/m01740_a0.cl index d4e4bbd64..84dc71ae0 100644 --- a/OpenCL/m01740_a0.cl +++ b/OpenCL/m01740_a0.cl @@ -30,7 +30,7 @@ __kernel void m01740_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha512_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m01740_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m01740_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha512_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m01740_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01750_a0.cl b/OpenCL/m01750_a0.cl index 65535a1ff..0f8743f85 100644 --- a/OpenCL/m01750_a0.cl +++ b/OpenCL/m01750_a0.cl @@ -30,7 +30,7 @@ __kernel void m01750_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m01750_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -95,7 +95,7 @@ __kernel void m01750_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -112,7 +112,7 @@ __kernel void m01750_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m01760_a0.cl b/OpenCL/m01760_a0.cl index d0ec8224b..0cd33c28a 100644 --- a/OpenCL/m01760_a0.cl +++ b/OpenCL/m01760_a0.cl @@ -30,7 +30,7 @@ __kernel void m01760_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -51,7 +51,7 @@ __kernel void m01760_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m01760_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -118,7 +118,7 @@ __kernel void m01760_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m02610_a0.cl b/OpenCL/m02610_a0.cl index 514af8e13..812dcb13d 100644 --- a/OpenCL/m02610_a0.cl +++ b/OpenCL/m02610_a0.cl @@ -60,7 +60,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -77,7 +77,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -175,7 +175,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -192,7 +192,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m02810_a0.cl b/OpenCL/m02810_a0.cl index 51c5185c6..4d488a79e 100644 --- a/OpenCL/m02810_a0.cl +++ b/OpenCL/m02810_a0.cl @@ -60,7 +60,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = 32; @@ -77,7 +77,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -198,7 +198,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = 32; @@ -215,7 +215,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m03710_a0.cl b/OpenCL/m03710_a0.cl index 89a878d63..9525c1587 100644 --- a/OpenCL/m03710_a0.cl +++ b/OpenCL/m03710_a0.cl @@ -60,7 +60,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -77,7 +77,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -188,7 +188,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -205,7 +205,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m03800_a0.cl b/OpenCL/m03800_a0.cl index b32882323..beba21325 100644 --- a/OpenCL/m03800_a0.cl +++ b/OpenCL/m03800_a0.cl @@ -30,7 +30,7 @@ __kernel void m03800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -53,7 +53,7 @@ __kernel void m03800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -101,7 +101,7 @@ __kernel void m03800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -124,7 +124,7 @@ __kernel void m03800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m03910_a0.cl b/OpenCL/m03910_a0.cl index 01a550ec3..86267ed95 100644 --- a/OpenCL/m03910_a0.cl +++ b/OpenCL/m03910_a0.cl @@ -60,7 +60,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = 32; @@ -77,7 +77,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -198,7 +198,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = 32; @@ -215,7 +215,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04010_a0.cl b/OpenCL/m04010_a0.cl index aba4c9cd3..239becda9 100644 --- a/OpenCL/m04010_a0.cl +++ b/OpenCL/m04010_a0.cl @@ -60,7 +60,7 @@ __kernel void m04010_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -74,7 +74,7 @@ __kernel void m04010_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -179,7 +179,7 @@ __kernel void m04010_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -193,7 +193,7 @@ __kernel void m04010_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04110_a0.cl b/OpenCL/m04110_a0.cl index 176b04d0c..4ce0d0f48 100644 --- a/OpenCL/m04110_a0.cl +++ b/OpenCL/m04110_a0.cl @@ -60,7 +60,7 @@ __kernel void m04110_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -83,7 +83,7 @@ __kernel void m04110_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -192,7 +192,7 @@ __kernel void m04110_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -215,7 +215,7 @@ __kernel void m04110_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04310_a0.cl b/OpenCL/m04310_a0.cl index 2f2e67619..5d94733bc 100644 --- a/OpenCL/m04310_a0.cl +++ b/OpenCL/m04310_a0.cl @@ -60,7 +60,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -77,7 +77,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -175,7 +175,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -192,7 +192,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04400_a0.cl b/OpenCL/m04400_a0.cl index 1a05232df..b62bb48a4 100644 --- a/OpenCL/m04400_a0.cl +++ b/OpenCL/m04400_a0.cl @@ -61,7 +61,7 @@ __kernel void m04400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -69,7 +69,7 @@ __kernel void m04400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -170,7 +170,7 @@ __kernel void m04400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -178,7 +178,7 @@ __kernel void m04400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04500_a0.cl b/OpenCL/m04500_a0.cl index fbb19bebc..0d8fc274c 100644 --- a/OpenCL/m04500_a0.cl +++ b/OpenCL/m04500_a0.cl @@ -60,7 +60,7 @@ __kernel void m04500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -68,7 +68,7 @@ __kernel void m04500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -169,7 +169,7 @@ __kernel void m04500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -177,7 +177,7 @@ __kernel void m04500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04520_a0.cl b/OpenCL/m04520_a0.cl index e3474f89f..0375fbd94 100644 --- a/OpenCL/m04520_a0.cl +++ b/OpenCL/m04520_a0.cl @@ -60,7 +60,7 @@ __kernel void m04520_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -74,7 +74,7 @@ __kernel void m04520_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -184,7 +184,7 @@ __kernel void m04520_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -198,7 +198,7 @@ __kernel void m04520_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04700_a0.cl b/OpenCL/m04700_a0.cl index 61524101a..585b870c0 100644 --- a/OpenCL/m04700_a0.cl +++ b/OpenCL/m04700_a0.cl @@ -61,7 +61,7 @@ __kernel void m04700_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -69,7 +69,7 @@ __kernel void m04700_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -166,7 +166,7 @@ __kernel void m04700_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -174,7 +174,7 @@ __kernel void m04700_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04800_a0.cl b/OpenCL/m04800_a0.cl index 614f6fdf0..c9053cc92 100644 --- a/OpenCL/m04800_a0.cl +++ b/OpenCL/m04800_a0.cl @@ -30,7 +30,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len - 1; @@ -55,7 +55,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -103,7 +103,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len - 1; @@ -128,7 +128,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m04900_a0.cl b/OpenCL/m04900_a0.cl index 3cf40686a..37f301775 100644 --- a/OpenCL/m04900_a0.cl +++ b/OpenCL/m04900_a0.cl @@ -30,7 +30,7 @@ __kernel void m04900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -53,7 +53,7 @@ __kernel void m04900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -101,7 +101,7 @@ __kernel void m04900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -124,7 +124,7 @@ __kernel void m04900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m05100_a0.cl b/OpenCL/m05100_a0.cl index 9900820b8..e8069da51 100644 --- a/OpenCL/m05100_a0.cl +++ b/OpenCL/m05100_a0.cl @@ -30,7 +30,7 @@ __kernel void m05100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m05100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m05100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -98,7 +98,7 @@ __kernel void m05100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m05300_a0.cl b/OpenCL/m05300_a0.cl index 8c52556d8..149443cc4 100644 --- a/OpenCL/m05300_a0.cl +++ b/OpenCL/m05300_a0.cl @@ -30,7 +30,7 @@ __kernel void m05300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m05300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -116,7 +116,7 @@ __kernel void m05300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -124,7 +124,7 @@ __kernel void m05300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m05400_a0.cl b/OpenCL/m05400_a0.cl index d89f33e82..3e560d501 100644 --- a/OpenCL/m05400_a0.cl +++ b/OpenCL/m05400_a0.cl @@ -30,7 +30,7 @@ __kernel void m05400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m05400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -116,7 +116,7 @@ __kernel void m05400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -124,7 +124,7 @@ __kernel void m05400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl index c11cff990..b9659b0f1 100644 --- a/OpenCL/m05500_a0.cl +++ b/OpenCL/m05500_a0.cl @@ -543,7 +543,7 @@ __kernel void m05500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -551,7 +551,7 @@ __kernel void m05500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -679,7 +679,7 @@ __kernel void m05500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -687,7 +687,7 @@ __kernel void m05500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m05600_a0.cl b/OpenCL/m05600_a0.cl index 722b475ca..c8d90140f 100644 --- a/OpenCL/m05600_a0.cl +++ b/OpenCL/m05600_a0.cl @@ -31,7 +31,7 @@ __kernel void m05600_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -39,7 +39,7 @@ __kernel void m05600_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -142,7 +142,7 @@ __kernel void m05600_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -150,7 +150,7 @@ __kernel void m05600_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m06000_a0.cl b/OpenCL/m06000_a0.cl index e9a367ec9..a59010df6 100644 --- a/OpenCL/m06000_a0.cl +++ b/OpenCL/m06000_a0.cl @@ -30,7 +30,7 @@ __kernel void m06000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m06000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m06000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m06000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m06100_a0.cl b/OpenCL/m06100_a0.cl index ccb3311dd..8522ff117 100644 --- a/OpenCL/m06100_a0.cl +++ b/OpenCL/m06100_a0.cl @@ -76,7 +76,7 @@ __kernel void m06100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -170,7 +170,7 @@ __kernel void m06100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m07000_a0.cl b/OpenCL/m07000_a0.cl index ae8beba28..cdef7f133 100644 --- a/OpenCL/m07000_a0.cl +++ b/OpenCL/m07000_a0.cl @@ -30,7 +30,7 @@ __kernel void m07000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m07000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -118,7 +118,7 @@ __kernel void m07000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -132,7 +132,7 @@ __kernel void m07000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m07300_a0.cl b/OpenCL/m07300_a0.cl index c552f3b02..62590732f 100644 --- a/OpenCL/m07300_a0.cl +++ b/OpenCL/m07300_a0.cl @@ -30,7 +30,7 @@ __kernel void m07300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m07300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m07300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m07300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0.cl index 9d6962738..a1a6f49b5 100644 --- a/OpenCL/m07500_a0.cl +++ b/OpenCL/m07500_a0.cl @@ -283,7 +283,7 @@ __kernel void m07500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); __local RC4_KEY rc4_keys[64]; @@ -311,7 +311,7 @@ __kernel void m07500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -352,7 +352,7 @@ __kernel void m07500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); __local RC4_KEY rc4_keys[64]; @@ -380,7 +380,7 @@ __kernel void m07500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m08100_a0.cl b/OpenCL/m08100_a0.cl index bfe18dddd..2efbff01c 100644 --- a/OpenCL/m08100_a0.cl +++ b/OpenCL/m08100_a0.cl @@ -30,7 +30,7 @@ __kernel void m08100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m08100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m08100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m08100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m08300_a0.cl b/OpenCL/m08300_a0.cl index 76d87c32d..9f4eae635 100644 --- a/OpenCL/m08300_a0.cl +++ b/OpenCL/m08300_a0.cl @@ -30,7 +30,7 @@ __kernel void m08300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -58,7 +58,7 @@ __kernel void m08300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -149,7 +149,7 @@ __kernel void m08300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -177,7 +177,7 @@ __kernel void m08300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m08400_a0.cl b/OpenCL/m08400_a0.cl index 402746bdd..fe4c4ed27 100644 --- a/OpenCL/m08400_a0.cl +++ b/OpenCL/m08400_a0.cl @@ -60,7 +60,7 @@ __kernel void m08400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -74,7 +74,7 @@ __kernel void m08400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -223,7 +223,7 @@ __kernel void m08400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -237,7 +237,7 @@ __kernel void m08400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m09900_a0.cl b/OpenCL/m09900_a0.cl index b0d3900ae..2701ca35e 100644 --- a/OpenCL/m09900_a0.cl +++ b/OpenCL/m09900_a0.cl @@ -30,7 +30,7 @@ __kernel void m09900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m09900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -86,7 +86,7 @@ __kernel void m09900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -94,7 +94,7 @@ __kernel void m09900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m10800_a0.cl b/OpenCL/m10800_a0.cl index f3185bda0..abe6efa67 100644 --- a/OpenCL/m10800_a0.cl +++ b/OpenCL/m10800_a0.cl @@ -45,7 +45,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -108,7 +108,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m11000_a0.cl b/OpenCL/m11000_a0.cl index 59adff088..8eec20c76 100644 --- a/OpenCL/m11000_a0.cl +++ b/OpenCL/m11000_a0.cl @@ -30,7 +30,7 @@ __kernel void m11000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m11000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -90,7 +90,7 @@ __kernel void m11000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -104,7 +104,7 @@ __kernel void m11000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m11100_a0.cl b/OpenCL/m11100_a0.cl index 6408a299c..a9d7e0f3a 100644 --- a/OpenCL/m11100_a0.cl +++ b/OpenCL/m11100_a0.cl @@ -82,7 +82,7 @@ __kernel void m11100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -90,7 +90,7 @@ __kernel void m11100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -240,7 +240,7 @@ __kernel void m11100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -248,7 +248,7 @@ __kernel void m11100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m11200_a0.cl b/OpenCL/m11200_a0.cl index b9d709a20..ca99cdb9c 100644 --- a/OpenCL/m11200_a0.cl +++ b/OpenCL/m11200_a0.cl @@ -30,7 +30,7 @@ __kernel void m11200_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -44,7 +44,7 @@ __kernel void m11200_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -158,7 +158,7 @@ __kernel void m11200_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -172,7 +172,7 @@ __kernel void m11200_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m11400_a0.cl b/OpenCL/m11400_a0.cl index 385b4484b..36408c7d3 100644 --- a/OpenCL/m11400_a0.cl +++ b/OpenCL/m11400_a0.cl @@ -60,7 +60,7 @@ __kernel void m11400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -74,7 +74,7 @@ __kernel void m11400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -170,7 +170,7 @@ __kernel void m11400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); md5_ctx_t ctx0; @@ -184,7 +184,7 @@ __kernel void m11400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m12600_a0.cl b/OpenCL/m12600_a0.cl index f7962bcef..a40bced95 100644 --- a/OpenCL/m12600_a0.cl +++ b/OpenCL/m12600_a0.cl @@ -76,7 +76,7 @@ __kernel void m12600_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -84,7 +84,7 @@ __kernel void m12600_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -224,7 +224,7 @@ __kernel void m12600_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -232,7 +232,7 @@ __kernel void m12600_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index 610dec8f3..2fe1b8399 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -392,7 +392,7 @@ __kernel void m13100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); __local RC4_KEY rc4_keys[64]; @@ -409,7 +409,7 @@ __kernel void m13100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -452,7 +452,7 @@ __kernel void m13100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); __local RC4_KEY rc4_keys[64]; @@ -469,7 +469,7 @@ __kernel void m13100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m13300_a0.cl b/OpenCL/m13300_a0.cl index 8b30c679f..52063f904 100644 --- a/OpenCL/m13300_a0.cl +++ b/OpenCL/m13300_a0.cl @@ -30,7 +30,7 @@ __kernel void m13300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m13300_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -88,7 +88,7 @@ __kernel void m13300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -96,7 +96,7 @@ __kernel void m13300_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m13500_a0.cl b/OpenCL/m13500_a0.cl index 6688bd602..33206b5d8 100644 --- a/OpenCL/m13500_a0.cl +++ b/OpenCL/m13500_a0.cl @@ -63,7 +63,7 @@ __kernel void m13500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -71,7 +71,7 @@ __kernel void m13500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -150,7 +150,7 @@ __kernel void m13500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -158,7 +158,7 @@ __kernel void m13500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m13800_a0.cl b/OpenCL/m13800_a0.cl index 3570447f5..e81a8afd4 100644 --- a/OpenCL/m13800_a0.cl +++ b/OpenCL/m13800_a0.cl @@ -30,7 +30,7 @@ __kernel void m13800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -38,7 +38,7 @@ __kernel void m13800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -88,7 +88,7 @@ __kernel void m13800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); /** * loop @@ -96,7 +96,7 @@ __kernel void m13800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m13900_a0.cl b/OpenCL/m13900_a0.cl index adf3e5ddf..c69c5be8b 100644 --- a/OpenCL/m13900_a0.cl +++ b/OpenCL/m13900_a0.cl @@ -60,7 +60,7 @@ __kernel void m13900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -74,7 +74,7 @@ __kernel void m13900_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -223,7 +223,7 @@ __kernel void m13900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -237,7 +237,7 @@ __kernel void m13900_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m14400_a0.cl b/OpenCL/m14400_a0.cl index aee068e19..cda06eb8b 100644 --- a/OpenCL/m14400_a0.cl +++ b/OpenCL/m14400_a0.cl @@ -60,7 +60,7 @@ __kernel void m14400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -122,7 +122,7 @@ __kernel void m14400_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -319,7 +319,7 @@ __kernel void m14400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); sha1_ctx_t ctx0; @@ -381,7 +381,7 @@ __kernel void m14400_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m15000_a0.cl b/OpenCL/m15000_a0.cl index 12f33300e..c49307c85 100644 --- a/OpenCL/m15000_a0.cl +++ b/OpenCL/m15000_a0.cl @@ -30,7 +30,7 @@ __kernel void m15000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m15000_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -97,7 +97,7 @@ __kernel void m15000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -114,7 +114,7 @@ __kernel void m15000_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); diff --git a/OpenCL/m15500_a0.cl b/OpenCL/m15500_a0.cl index 6013e9b3e..a0d36724e 100644 --- a/OpenCL/m15500_a0.cl +++ b/OpenCL/m15500_a0.cl @@ -30,7 +30,7 @@ __kernel void m15500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -47,7 +47,7 @@ __kernel void m15500_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len); @@ -103,7 +103,7 @@ __kernel void m15500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru * base */ - pw_t pw = pws[gid]; + COPY_PW (pws[gid]); const u32 salt_len = salt_bufs[salt_pos].salt_len; @@ -120,7 +120,7 @@ __kernel void m15500_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) { - pw_t tmp = pw; + pw_t tmp = PASTE_PW; tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);