Optimized a0 pure kernel for AMD

pull/1342/head^2
jsteube 7 years ago
parent a5659d5619
commit 0783289e2f

@ -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

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

@ -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);

Loading…
Cancel
Save