diff --git a/OpenCL/m12500.cl b/OpenCL/m12500.cl
index a8d4a5263..3290d96f5 100644
--- a/OpenCL/m12500.cl
+++ b/OpenCL/m12500.cl
@@ -1079,7 +1079,7 @@ __kernel void m12500_loop (__global pw_t *pws, __global const kernel_rule_t *rul
         iter++;
       }
 
-      for (u32 j = 0; j < p3; j++)
+      for (volatile u32 j = 0; j < p3; j++)
       {
         const u32 j16 = j * 16;
 
diff --git a/docs/changes.txt b/docs/changes.txt
index aa76d7d0d..2211c79e5 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -32,6 +32,7 @@
 ##
 
 - Workaround added for AMDGPU-Pro OpenCL runtime: AES encrypt and decrypt Invertkey function was calculated wrong in certain cases
+- Workaround added for AMDGPU-Pro OpenCL runtime: RAR3 kernel require a volatile variable to work correctly
 
 ##
 ## Bugs