From 1d43540fc4fae07af5b65b4f465bc905fc8431e8 Mon Sep 17 00:00:00 2001 From: unix-ninja Date: Wed, 17 Oct 2018 11:03:20 -0400 Subject: [PATCH] Simplify alignment masks for mode 18100 --- OpenCL/m18100_a0-pure.cl | 12 ++++++------ OpenCL/m18100_a3-pure.cl | 14 +++++++------- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/OpenCL/m18100_a0-pure.cl b/OpenCL/m18100_a0-pure.cl index 76a7f2705..7ab09d5d1 100644 --- a/OpenCL/m18100_a0-pure.cl +++ b/OpenCL/m18100_a0-pure.cl @@ -70,13 +70,13 @@ __kernel void m18100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; @@ -163,13 +163,13 @@ __kernel void m18100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; diff --git a/OpenCL/m18100_a3-pure.cl b/OpenCL/m18100_a3-pure.cl index 8f90361e9..1352338a1 100644 --- a/OpenCL/m18100_a3-pure.cl +++ b/OpenCL/m18100_a3-pure.cl @@ -79,13 +79,13 @@ __kernel void m18100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule switch(otp_offset&3) { case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfeffffff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4]; @@ -182,14 +182,14 @@ __kernel void m18100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule // on some systems, &3 is faster than %4, so we will use it in our switch() switch(otp_offset&3) { - case 1: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 24); + case 1: + otp_code = ((ctx.opad.h[otp_offset/4] & 0x00ffffff) << 8) | ((ctx.opad.h[otp_offset/4+1] & 0xff000000) >> 24); break; case 2: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xfffeffff) >> 16); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x0000ffff) << 16) | ((ctx.opad.h[otp_offset/4+1] & 0xffff0000) >> 16); break; case 3: - otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xfffffeff) >> 8); + otp_code = ((ctx.opad.h[otp_offset/4] & 0x000000ff) << 24) | ((ctx.opad.h[otp_offset/4+1] & 0xffffff00) >> 8); break; default: otp_code = ctx.opad.h[otp_offset/4];