diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl
index 8fccfd6b5..e9b845ea3 100644
--- a/OpenCL/m08500_a0.cl
+++ b/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);
 
diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl
index 63ee260e6..65bb133ee 100644
--- a/OpenCL/m08500_a1.cl
+++ b/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);
 
diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3.cl
index 933261a48..dba7ada15 100644
--- a/OpenCL/m08500_a3.cl
+++ b/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);