Fix -m 8500: volatile no longer required

OpenCL/m08500_a0.cl

@@ -631,7 +631,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];
 
-    volatile u32x iv[2];
+    u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 
@@ -764,7 +764,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];
 
-    volatile u32x iv[2];
+    u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 

OpenCL/m08500_a1.cl

@@ -671,7 +671,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];
 
-    volatile u32x iv[2];
+    u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 
@@ -847,7 +847,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];
 
-    volatile u32x iv[2];
+    u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 

OpenCL/m08500_a3.cl

@@ -579,7 +579,7 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    volatile u32x iv[2];
+    u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
 
@@ -654,7 +654,7 @@ void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
     data[0] = salt_buf0[0];
     data[1] = salt_buf0[1];
 
-    volatile u32x iv[2];
+    u32x iv[2];
 
     _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);