From f3a394b85da5f3efdd741037adc6d15e0180125c Mon Sep 17 00:00:00 2001 From: jsteube Date: Thu, 24 Dec 2015 12:10:03 +0100 Subject: [PATCH] Workaround buggy NV OpenCL runtime --- OpenCL/m05500_a0.cl | 5 ++++- OpenCL/m05500_a1.cl | 5 ++++- OpenCL/m05500_a3.cl | 5 ++++- 3 files changed, 12 insertions(+), 3 deletions(-) diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl index 64d1a40..456a656 100644 --- a/OpenCL/m05500_a0.cl +++ b/OpenCL/m05500_a0.cl @@ -673,7 +673,10 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo * DES2 */ - transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); + volatile const u32 bc = (b >> 24) | (c << 8); + volatile const u32 cd = (c >> 24) | (d << 8); + + transform_netntlmv1_key (bc, cd, key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); diff --git a/OpenCL/m05500_a1.cl b/OpenCL/m05500_a1.cl index 4d44506..59f5277 100644 --- a/OpenCL/m05500_a1.cl +++ b/OpenCL/m05500_a1.cl @@ -725,7 +725,10 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_m04 (__glo * DES2 */ - transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key); + volatile const u32 bc = (b >> 24) | (c << 8); + volatile const u32 cd = (c >> 24) | (d << 8); + + transform_netntlmv1_key (bc, cd, key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); diff --git a/OpenCL/m05500_a3.cl b/OpenCL/m05500_a3.cl index 3a507f4..b798d14 100644 --- a/OpenCL/m05500_a3.cl +++ b/OpenCL/m05500_a3.cl @@ -602,7 +602,10 @@ 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); + volatile const u32 bc = (b >> 24) | (c << 8); + volatile const u32 cd = (c >> 24) | (d << 8); + + transform_netntlmv1_key (bc, cd, key); _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb); -- 2.43.0