Workaround buggy NV OpenCL runtime

diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl
index 64d1a40..456a656 100644 (file)
--- 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 (file)
--- 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 (file)
--- 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);