From 0522ca9cbedc7b30ad358df0cf2cf69695f1d48f Mon Sep 17 00:00:00 2001 From: jsteube Date: Thu, 24 Dec 2015 12:05:43 +0100 Subject: [PATCH] Revert "Workaround buggy NV OpenCL runtime" This reverts commit 90eba9cd2e1994b6eb01a858e4603936e2c50e83. --- OpenCL/m05500_a0.cl | 17 +++++++++-------- OpenCL/m05500_a1.cl | 17 +++++++++-------- OpenCL/m05500_a3.cl | 19 ++++++++----------- 3 files changed, 26 insertions(+), 27 deletions(-) diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl index fce3625f6..64d1a40a9 100644 --- a/OpenCL/m05500_a0.cl +++ b/OpenCL/m05500_a0.cl @@ -673,10 +673,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo * DES2 */ - volatile const u32 bc = (b >> 24) | (c << 8); - volatile const u32 cd = (c >> 24) | (d << 8); - - transform_netntlmv1_key (bc, cd, key); + transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); @@ -922,16 +919,20 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__glo * DES2 */ - volatile const u32 bc = (b >> 24) | (c << 8); - volatile const u32 cd = (c >> 24) | (d << 8); - - transform_netntlmv1_key (bc, cd, key); + /* + transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); u32 iv2[2]; _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); + */ + + u32 iv2[2]; + + iv2[0] = search[2]; + iv2[1] = search[3]; /** * compare diff --git a/OpenCL/m05500_a1.cl b/OpenCL/m05500_a1.cl index 81d76503c..4d44506b0 100644 --- a/OpenCL/m05500_a1.cl +++ b/OpenCL/m05500_a1.cl @@ -725,10 +725,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo * DES2 */ - volatile const u32 bc = (b >> 24) | (c << 8); - volatile const u32 cd = (c >> 24) | (d << 8); - - transform_netntlmv1_key (bc, cd, key); + transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); @@ -1028,16 +1025,20 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__glo * DES2 */ - volatile const u32 bc = (b >> 24) | (c << 8); - volatile const u32 cd = (c >> 24) | (d << 8); - - transform_netntlmv1_key (bc, cd, key); + /* + transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); u32 iv2[2]; _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); + */ + + u32 iv2[2]; + + iv2[0] = search[2]; + iv2[1] = search[3]; /** * compare diff --git a/OpenCL/m05500_a3.cl b/OpenCL/m05500_a3.cl index 4b08e140c..3a507f453 100644 --- a/OpenCL/m05500_a3.cl +++ b/OpenCL/m05500_a3.cl @@ -602,7 +602,6 @@ static void m05500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 * DES2 */ - /* transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); @@ -610,12 +609,6 @@ static void m05500m (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 u32 iv2[2]; _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); - */ - - u32 iv2[2]; - - iv2[0] = search[2]; - iv2[1] = search[3]; /** * compare @@ -778,16 +771,20 @@ static void m05500s (__local u32 s_SPtrans[8][64], __local u32 s_skb[8][64], u32 * DES2 */ - volatile const u32 bc = (b >> 24) | (c << 8); - volatile const u32 cd = (c >> 24) | (d << 8); - - transform_netntlmv1_key (bc, cd, key); + /* + transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); u32 iv2[2]; _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans); + */ + + u32 iv2[2]; + + iv2[0] = search[2]; + iv2[1] = search[3]; /** * compare