Workaround buggy NV OpenCL runtime

diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl
index 64d1a40..fce3625 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);
 
@@ -919,20 +922,16 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__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);
 
     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 4d44506..81d7650 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);
 
@@ -1025,20 +1028,16 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05500_s04 (__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);
 
     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 3a507f4..4b08e14 100644 (file)
--- a/OpenCL/m05500_a3.cl
+++ b/OpenCL/m05500_a3.cl
@@ -602,6 +602,7 @@ 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);
@@ -609,6 +610,12 @@ 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
@@ -771,20 +778,16 @@ static void m05500s (__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);
 
     u32 iv2[2];
 
     _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
-    */
-
-    u32 iv2[2];
-
-    iv2[0] = search[2];
-    iv2[1] = search[3];
 
     /**
      * compare