github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-11-02 20:39:22 +01:00
Code Releases Activity

Remove global barrier when not needed to workaround Intel OpenCL runtime bug

Browse Source
This commit is contained in:
jsteube 2017-08-07 17:25:15 +02:00
parent b9876c100b
commit 0a676b549f
180 changed files with 0 additions and 978 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
OpenCL/m00000_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00000_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00010_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00010_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m00010_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -123,8 +119,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00020_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

4
OpenCL/m00020_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m00030_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00030_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m00030_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -123,8 +119,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00040_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

4
OpenCL/m00040_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m00050_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -117,8 +113,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -130,8 +124,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00050_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -135,8 +131,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -148,8 +142,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00050_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -121,8 +117,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -134,8 +128,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00060_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_hmac_ctx_t ctx0; md5_hmac_ctx_t ctx0;
@ -119,8 +115,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_hmac_ctx_t ctx0; md5_hmac_ctx_t ctx0;

8
OpenCL/m00060_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_hmac_ctx_t ctx0; md5_hmac_ctx_t ctx0;
@ -137,8 +133,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -150,8 +144,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_hmac_ctx_t ctx0; md5_hmac_ctx_t ctx0;

8
OpenCL/m00060_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_hmac_ctx_vector_t ctx0; md5_hmac_ctx_vector_t ctx0;
@ -123,8 +119,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_hmac_ctx_vector_t ctx0; md5_hmac_ctx_vector_t ctx0;

4
OpenCL/m00100_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00100_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00110_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00110_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;

8
OpenCL/m00110_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -123,8 +119,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00120_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;

4
OpenCL/m00120_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;

8
OpenCL/m00130_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00130_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;

8
OpenCL/m00130_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -124,8 +120,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -137,8 +131,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00140_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;

4
OpenCL/m00140_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_ctx_t ctx0; sha1_ctx_t ctx0;

8
OpenCL/m00150_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -117,8 +113,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -130,8 +124,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00150_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -135,8 +131,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -148,8 +142,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00150_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -121,8 +117,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -134,8 +128,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m00160_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_hmac_ctx_t ctx0; sha1_hmac_ctx_t ctx0;
@ -119,8 +115,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_hmac_ctx_t ctx0; sha1_hmac_ctx_t ctx0;

8
OpenCL/m00160_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_hmac_ctx_t ctx0; sha1_hmac_ctx_t ctx0;
@ -137,8 +133,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -150,8 +144,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_hmac_ctx_t ctx0; sha1_hmac_ctx_t ctx0;

8
OpenCL/m00160_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_hmac_ctx_vector_t ctx0; sha1_hmac_ctx_vector_t ctx0;
@ -123,8 +119,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha1_hmac_ctx_vector_t ctx0; sha1_hmac_ctx_vector_t ctx0;

4
OpenCL/m00300_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -127,8 +125,6 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00300_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -131,8 +129,6 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

2
OpenCL/m00400.cl
View File

@ -76,8 +76,6 @@ __kernel void m00400_loop (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
u32 digest[4]; u32 digest[4];

8
OpenCL/m00500.cl
View File

@ -40,8 +40,6 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -53,8 +51,6 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -154,8 +150,6 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -167,8 +161,6 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00900_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m00900_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01000_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01000_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01100_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -132,8 +128,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -145,8 +139,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01100_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md4_ctx_t ctx0; md4_ctx_t ctx0;
@ -119,8 +117,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md4_ctx_t ctx0; md4_ctx_t ctx0;

8
OpenCL/m01100_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -136,8 +132,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -149,8 +143,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01300_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01300_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01400_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01400_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01410_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01410_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;

8
OpenCL/m01410_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -123,8 +119,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01420_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;

4
OpenCL/m01420_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;

8
OpenCL/m01430_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01430_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;

8
OpenCL/m01430_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -124,8 +120,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -137,8 +131,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01440_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;

4
OpenCL/m01440_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_ctx_t ctx0; sha256_ctx_t ctx0;

8
OpenCL/m01450_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -117,8 +113,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -130,8 +124,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01450_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -135,8 +131,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -148,8 +142,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01450_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -121,8 +117,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -134,8 +128,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01460_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_hmac_ctx_t ctx0; sha256_hmac_ctx_t ctx0;
@ -119,8 +115,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_hmac_ctx_t ctx0; sha256_hmac_ctx_t ctx0;

8
OpenCL/m01460_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_hmac_ctx_t ctx0; sha256_hmac_ctx_t ctx0;
@ -137,8 +133,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -150,8 +144,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_hmac_ctx_t ctx0; sha256_hmac_ctx_t ctx0;

8
OpenCL/m01460_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_hmac_ctx_vector_t ctx0; sha256_hmac_ctx_vector_t ctx0;
@ -123,8 +119,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha256_hmac_ctx_vector_t ctx0; sha256_hmac_ctx_vector_t ctx0;

8
OpenCL/m01600.cl
View File

@ -41,8 +41,6 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -54,8 +52,6 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -156,8 +152,6 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -169,8 +163,6 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01700_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -104,8 +102,6 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01700_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -108,8 +106,6 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01710_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01710_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;

8
OpenCL/m01710_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -123,8 +119,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01720_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;

4
OpenCL/m01720_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;

8
OpenCL/m01730_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -119,8 +115,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01730_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;
@ -106,8 +104,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;

8
OpenCL/m01730_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -124,8 +120,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -137,8 +131,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m01740_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;

4
OpenCL/m01740_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;
@ -114,8 +112,6 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_ctx_t ctx0; sha512_ctx_t ctx0;

8
OpenCL/m01750_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -117,8 +113,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -130,8 +124,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01750_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -135,8 +131,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -148,8 +142,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01750_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -121,8 +117,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -134,8 +128,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m01760_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_hmac_ctx_t ctx0; sha512_hmac_ctx_t ctx0;
@ -119,8 +115,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -132,8 +126,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_hmac_ctx_t ctx0; sha512_hmac_ctx_t ctx0;

8
OpenCL/m01760_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_hmac_ctx_t ctx0; sha512_hmac_ctx_t ctx0;
@ -137,8 +133,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = swap32_S (pws[gid].i[idx]); w[idx] = swap32_S (pws[gid].i[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -150,8 +144,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_hmac_ctx_t ctx0; sha512_hmac_ctx_t ctx0;

8
OpenCL/m01760_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_hmac_ctx_vector_t ctx0; sha512_hmac_ctx_vector_t ctx0;
@ -123,8 +119,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]); s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
barrier (CLK_GLOBAL_MEM_FENCE);
} }
sha512_hmac_ctx_vector_t ctx0; sha512_hmac_ctx_vector_t ctx0;

4
OpenCL/m01800.cl
View File

@ -36,8 +36,6 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
@ -54,8 +52,6 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)

8
OpenCL/m02610_a0.cl
View File

@ -69,8 +69,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -82,8 +80,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -197,8 +193,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -210,8 +204,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m02610_a1.cl
View File

@ -67,8 +67,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -184,8 +182,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m02610_a3.cl
View File

@ -67,8 +67,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -80,8 +78,6 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -201,8 +197,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -214,8 +208,6 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m02810_a0.cl
View File

@ -69,8 +69,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = 32; const u32 salt_len = 32;
@ -82,8 +80,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -220,8 +216,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = 32; const u32 salt_len = 32;
@ -233,8 +227,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m02810_a1.cl
View File

@ -67,8 +67,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -207,8 +205,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m02810_a3.cl
View File

@ -67,8 +67,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = 32; const u32 salt_len = 32;
@ -80,8 +78,6 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -224,8 +220,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = 32; const u32 salt_len = 32;
@ -237,8 +231,6 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m03710_a0.cl
View File

@ -69,8 +69,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -82,8 +80,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -210,8 +206,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -223,8 +217,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

4
OpenCL/m03710_a1.cl
View File

@ -67,8 +67,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -197,8 +195,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m03710_a3.cl
View File

@ -67,8 +67,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -80,8 +78,6 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -214,8 +210,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -227,8 +221,6 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

8
OpenCL/m03800_a0.cl
View File

@ -39,8 +39,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -52,8 +50,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -123,8 +119,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -136,8 +130,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

4
OpenCL/m03800_a1.cl
View File

@ -37,8 +37,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -108,8 +106,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m03800_a3.cl
View File

@ -37,8 +37,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -50,8 +48,6 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;
@ -129,8 +125,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = salt_bufs[salt_pos].salt_len; const u32 salt_len = salt_bufs[salt_pos].salt_len;
@ -142,8 +136,6 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf[idx]; s[idx] = salt_bufs[salt_pos].salt_buf[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
md5_ctx_t ctx0; md5_ctx_t ctx0;

8
OpenCL/m03910_a0.cl
View File

@ -69,8 +69,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = 32; const u32 salt_len = 32;
@ -82,8 +80,6 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**
@ -220,8 +216,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < pw_lenv; idx++) for (int idx = 0; idx < pw_lenv; idx++)
{ {
w[idx] = pws[gid].i[idx]; w[idx] = pws[gid].i[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
const u32 salt_len = 32; const u32 salt_len = 32;
@ -233,8 +227,6 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
for (int idx = 0; idx < salt_lenv; idx++) for (int idx = 0; idx < salt_lenv; idx++)
{ {
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx]; s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
barrier (CLK_GLOBAL_MEM_FENCE);
} }
/** /**

Some files were not shown because too many files have changed in this diff Show More