diff --git a/OpenCL/m01700_a0-pure.cl b/OpenCL/m01700_a0-pure.cl index c1b52cf10..31ac12008 100644 --- a/OpenCL/m01700_a0-pure.cl +++ b/OpenCL/m01700_a0-pure.cl @@ -38,7 +38,7 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { - w[idx] = pws[gid].i[idx]; + w[idx] = swap32_S (pws[gid].i[idx]); barrier (CLK_GLOBAL_MEM_FENCE); } @@ -103,7 +103,7 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { - w[idx] = pws[gid].i[idx]; + w[idx] = swap32_S (pws[gid].i[idx]); barrier (CLK_GLOBAL_MEM_FENCE); } diff --git a/OpenCL/m01700_a1-pure.cl b/OpenCL/m01700_a1-pure.cl index f493ac05f..a0fc5df3a 100644 --- a/OpenCL/m01700_a1-pure.cl +++ b/OpenCL/m01700_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha512_init (&ctx0); - sha512_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); /** * loop @@ -42,7 +42,7 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule { sha512_ctx_t ctx = ctx0; - sha512_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); sha512_final (&ctx); @@ -86,7 +86,7 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha512_init (&ctx0); - sha512_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); /** * loop @@ -96,7 +96,7 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule { sha512_ctx_t ctx = ctx0; - sha512_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); sha512_final (&ctx); diff --git a/OpenCL/m10800_a0-pure.cl b/OpenCL/m10800_a0-pure.cl index c0f01faa6..1ae43d030 100644 --- a/OpenCL/m10800_a0-pure.cl +++ b/OpenCL/m10800_a0-pure.cl @@ -38,7 +38,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { - w[idx] = pws[gid].i[idx]; + w[idx] = swap32_S (pws[gid].i[idx]); barrier (CLK_GLOBAL_MEM_FENCE); } @@ -103,7 +103,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule for (int idx = 0; idx < pw_lenv; idx++) { - w[idx] = pws[gid].i[idx]; + w[idx] = swap32_S (pws[gid].i[idx]); barrier (CLK_GLOBAL_MEM_FENCE); } diff --git a/OpenCL/m10800_a1-pure.cl b/OpenCL/m10800_a1-pure.cl index 1fa07f4a8..9f1eded3a 100644 --- a/OpenCL/m10800_a1-pure.cl +++ b/OpenCL/m10800_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule sha384_init (&ctx0); - sha384_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); /** * loop @@ -42,7 +42,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule { sha384_ctx_t ctx = ctx0; - sha384_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); sha384_final (&ctx); @@ -86,7 +86,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule sha384_init (&ctx0); - sha384_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); /** * loop @@ -96,7 +96,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule { sha384_ctx_t ctx = ctx0; - sha384_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); sha384_final (&ctx);