Fix warp bug on -m 8500

diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl
index 15dc885..71bdc28 100644 (file)
--- a/OpenCL/m08500_a0.cl
+++ b/OpenCL/m08500_a0.cl
@@ -626,7 +626,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    u32x iv[2];
+    volatile u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 
@@ -760,7 +760,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    u32x iv[2];
+    volatile u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 

diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl
index ccefd47..f3f653d 100644 (file)
--- a/OpenCL/m08500_a1.cl
+++ b/OpenCL/m08500_a1.cl
@@ -685,7 +685,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    u32 iv[2];
+    volatile u32 iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 
@@ -898,7 +898,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    u32 iv[2];
+    volatile u32 iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 

diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3.cl
index ebc8470..2e9b9c5 100644 (file)
--- a/OpenCL/m08500_a3.cl
+++ b/OpenCL/m08500_a3.cl
@@ -570,7 +570,7 @@ static void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    u32x iv[2];
+    volatile u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 
@@ -642,7 +642,7 @@ static void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    u32x iv[2];
+    volatile u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);