Reformatted unrollable kernel loops and added the new promote_to_registers pragma for several kernels

7 files changed, 133 insertions, 133 deletions

diff --git a/src/kernel_preprocessor.cpp b/src/kernel_preprocessor.cpp
index b6c0b398..6aa853fd 100644
--- a/src/kernel_preprocessor.cpp
+++ b/src/kernel_preprocessor.cpp
@@ -164,9 +164,6 @@ std::vector<std::string> PreprocessDefinesAndComments(const std::string& source,
         if (HasOnlyDigits(value)) {
           defines.emplace(name, std::stoi(value));
         }
-        else {
-          printf("'%s'\n", value.c_str());
-        }
       }
 
       // Detect #ifndef blocks
diff --git a/src/kernels/level1/xaxpy.opencl b/src/kernels/level1/xaxpy.opencl
index 3a574ec2..74e49930 100644
--- a/src/kernels/level1/xaxpy.opencl
+++ b/src/kernels/level1/xaxpy.opencl
@@ -45,8 +45,8 @@ void XaxpyFaster(const int n, const real_arg arg_alpha,
 
   if (get_global_id(0) < n / (VW)) {
     #pragma unroll
-    for (int w = 0; w < WPT; w += 1) {
-      const int id = w*get_global_size(0) + get_global_id(0);
+    for (int _w = 0; _w < WPT; _w += 1) {
+      const int id = _w*get_global_size(0) + get_global_id(0);
       realV xvalue = xgm[id];
       realV yvalue = ygm[id];
       ygm[id] = MultiplyAddVector(yvalue, alpha, xvalue);
@@ -63,8 +63,8 @@ void XaxpyFastest(const int n, const real_arg arg_alpha,
   const real alpha = GetRealArg(arg_alpha);
 
   #pragma unroll
-  for (int w = 0; w < WPT; w += 1) {
-    const int id = w*get_global_size(0) + get_global_id(0);
+  for (int _w = 0; _w < WPT; _w += 1) {
+    const int id = _w*get_global_size(0) + get_global_id(0);
     realV xvalue = xgm[id];
     realV yvalue = ygm[id];
     ygm[id] = MultiplyAddVector(yvalue, alpha, xvalue);
diff --git a/src/kernels/level2/xgemv.opencl b/src/kernels/level2/xgemv.opencl
index ea0478f0..2a50e8fb 100644
--- a/src/kernels/level2/xgemv.opencl
+++ b/src/kernels/level2/xgemv.opencl
@@ -227,10 +227,11 @@ void Xgemv(const int m, const int n,
   __local real xlm[WGS1];
 
   // Initializes the accumulation register
+  #pragma promote_to_registers
   real acc[WPT1];
   #pragma unroll
-  for (int w=0; w<WPT1; ++w) {
-    SetToZero(acc[w]);
+  for (int _w = 0; _w < WPT1; _w += 1) {
+    SetToZero(acc[_w]);
   }
 
   // Divides the work in a main and tail section
@@ -248,30 +249,31 @@ void Xgemv(const int m, const int n,
     barrier(CLK_LOCAL_MEM_FENCE);
 
     // Loops over the work per thread, and checks whether in bounds
-    for (int w=0; w<WPT1; ++w) {
-      const int gid = w*get_global_size(0) + get_global_id(0);
+    #pragma unroll
+    for (int _w = 0; _w < WPT1; _w += 1) {
+      const int gid = _w*get_global_size(0) + get_global_id(0);
       if (gid < m) {
 
         // The multiply-add function for the main part (divisable by WGS1)
         if (a_rotated == 0) { // Not rotated
           for (int kloop=0; kloop<WGS1; kloop+=UNROLL1) {
             #pragma unroll
-            for (int kunroll=0; kunroll<UNROLL1; ++kunroll) {
-              const int k = kwg + kloop + kunroll;
+            for (int _kunroll = 0; _kunroll < UNROLL1; _kunroll += 1) {
+              const int k = kwg + kloop + _kunroll;
               real value = LoadMatrixA(agm, gid, k, a_ld, a_offset, parameter, kl, ku);
               if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
-              MultiplyAdd(acc[w], xlm[kloop + kunroll], value);
+              MultiplyAdd(acc[_w], xlm[kloop + _kunroll], value);
             }
           }
         }
         else { // Transposed
           for (int kloop=0; kloop<WGS1; kloop+=UNROLL1) {
             #pragma unroll
-            for (int kunroll=0; kunroll<UNROLL1; ++kunroll) {
-              const int k = kwg + kloop + kunroll;
+            for (int _kunroll = 0; _kunroll < UNROLL1; _kunroll += 1) {
+              const int k = kwg + kloop + _kunroll;
               real value = LoadMatrixA(agm, k, gid, a_ld, a_offset, parameter, kl, ku);
               if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
-              MultiplyAdd(acc[w], xlm[kloop + kunroll], value);
+              MultiplyAdd(acc[_w], xlm[kloop + _kunroll], value);
             }
           }
         }
@@ -284,31 +286,29 @@ void Xgemv(const int m, const int n,
 
   // Loops over the work per thread, and checks whether in bounds
   #pragma unroll
-  for (int w=0; w<WPT1; ++w) {
-    const int gid = w*get_global_size(0) + get_global_id(0);
+  for (int _w = 0; _w < WPT1; _w += 1) {
+    const int gid = _w*get_global_size(0) + get_global_id(0);
     if (gid < m) {
 
       // The multiply-add function for the remainder part (not divisable by WGS1)
       if (a_rotated == 0) { // Not rotated
-        #pragma unroll
         for (int k=n_floor; k<n; ++k) {
           real value = LoadMatrixA(agm, gid, k, a_ld, a_offset, parameter, kl, ku);
           if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
-          MultiplyAdd(acc[w], xgm[k*x_inc + x_offset], value);
+          MultiplyAdd(acc[_w], xgm[k*x_inc + x_offset], value);
         }
       }
       else { // Transposed
-        #pragma unroll
         for (int k=n_floor; k<n; ++k) {
           real value = LoadMatrixA(agm, k, gid, a_ld, a_offset, parameter, kl, ku);
           if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
-          MultiplyAdd(acc[w], xgm[k*x_inc + x_offset], value);
+          MultiplyAdd(acc[_w], xgm[k*x_inc + x_offset], value);
         }
       }
 
       // Stores the final result
       real yval = ygm[gid*y_inc + y_offset];
-      AXPBY(ygm[gid*y_inc + y_offset], alpha, acc[w], beta, yval);
+      AXPBY(ygm[gid*y_inc + y_offset], alpha, acc[_w], beta, yval);
     }
   }
 }
diff --git a/src/kernels/level2/xgemv_fast.opencl b/src/kernels/level2/xgemv_fast.opencl
index 8a08f076..892bc55c 100644
--- a/src/kernels/level2/xgemv_fast.opencl
+++ b/src/kernels/level2/xgemv_fast.opencl
@@ -105,10 +105,11 @@ void XgemvFast(const int m, const int n,
   __local real xlm[WGS2];
 
   // Initializes the accumulation registers
+  #pragma promote_to_registers
   real acc[WPT2];
   #pragma unroll
-  for (int w=0; w<WPT2; ++w) {
-    SetToZero(acc[w]);
+  for (int _w = 0; _w < WPT2; _w += 1) {
+    SetToZero(acc[_w]);
   }
 
   // Loops over work-group sized portions of the work
@@ -123,48 +124,48 @@ void XgemvFast(const int m, const int n,
 
     // The multiply-add function (not rotated)
     #pragma unroll
-    for (int kl=0; kl<WGS2; ++kl) {
-      const int k = kwg + kl;
+    for (int _kl = 0; _kl < WGS2; _kl += 1) {
+      const int k = kwg + _kl;
       #pragma unroll
-      for (int w=0; w<WPT2/VW2; ++w) {
-        const int gid = (WPT2/VW2)*get_global_id(0) + w;
+      for (int _w = 0; _w < WPT2/VW2; _w += 1) {
+        const int gid = (WPT2/VW2)*get_global_id(0) + _w;
         realVF avec = agm[(a_ld/VW2)*k + gid];
         #if VW2 == 1
-          MultiplyAdd(acc[VW2*w+0], xlm[kl], avec);
+          MultiplyAdd(acc[VW2*_w+0], xlm[_kl], avec);
         #elif VW2 == 2
-          MultiplyAdd(acc[VW2*w+0], xlm[kl], avec.x);
-          MultiplyAdd(acc[VW2*w+1], xlm[kl], avec.y);
+          MultiplyAdd(acc[VW2*_w+0], xlm[_kl], avec.x);
+          MultiplyAdd(acc[VW2*_w+1], xlm[_kl], avec.y);
         #elif VW2 == 4
-          MultiplyAdd(acc[VW2*w+0], xlm[kl], avec.x);
-          MultiplyAdd(acc[VW2*w+1], xlm[kl], avec.y);
-          MultiplyAdd(acc[VW2*w+2], xlm[kl], avec.z);
-          MultiplyAdd(acc[VW2*w+3], xlm[kl], avec.w);
+          MultiplyAdd(acc[VW2*_w+0], xlm[_kl], avec.x);
+          MultiplyAdd(acc[VW2*_w+1], xlm[_kl], avec.y);
+          MultiplyAdd(acc[VW2*_w+2], xlm[_kl], avec.z);
+          MultiplyAdd(acc[VW2*_w+3], xlm[_kl], avec.w);
         #elif VW2 == 8
-          MultiplyAdd(acc[VW2*w+0], xlm[kl], avec.s0);
-          MultiplyAdd(acc[VW2*w+1], xlm[kl], avec.s1);
-          MultiplyAdd(acc[VW2*w+2], xlm[kl], avec.s2);
-          MultiplyAdd(acc[VW2*w+3], xlm[kl], avec.s3);
-          MultiplyAdd(acc[VW2*w+4], xlm[kl], avec.s4);
-          MultiplyAdd(acc[VW2*w+5], xlm[kl], avec.s5);
-          MultiplyAdd(acc[VW2*w+6], xlm[kl], avec.s6);
-          MultiplyAdd(acc[VW2*w+7], xlm[kl], avec.s7);
+          MultiplyAdd(acc[VW2*_w+0], xlm[_kl], avec.s0);
+          MultiplyAdd(acc[VW2*_w+1], xlm[_kl], avec.s1);
+          MultiplyAdd(acc[VW2*_w+2], xlm[_kl], avec.s2);
+          MultiplyAdd(acc[VW2*_w+3], xlm[_kl], avec.s3);
+          MultiplyAdd(acc[VW2*_w+4], xlm[_kl], avec.s4);
+          MultiplyAdd(acc[VW2*_w+5], xlm[_kl], avec.s5);
+          MultiplyAdd(acc[VW2*_w+6], xlm[_kl], avec.s6);
+          MultiplyAdd(acc[VW2*_w+7], xlm[_kl], avec.s7);
         #elif VW2 == 16
-          MultiplyAdd(acc[VW2*w+0], xlm[kl], avec.s0);
-          MultiplyAdd(acc[VW2*w+1], xlm[kl], avec.s1);
-          MultiplyAdd(acc[VW2*w+2], xlm[kl], avec.s2);
-          MultiplyAdd(acc[VW2*w+3], xlm[kl], avec.s3);
-          MultiplyAdd(acc[VW2*w+4], xlm[kl], avec.s4);
-          MultiplyAdd(acc[VW2*w+5], xlm[kl], avec.s5);
-          MultiplyAdd(acc[VW2*w+6], xlm[kl], avec.s6);
-          MultiplyAdd(acc[VW2*w+7], xlm[kl], avec.s7);
-          MultiplyAdd(acc[VW2*w+8], xlm[kl], avec.s8);
-          MultiplyAdd(acc[VW2*w+9], xlm[kl], avec.s9);
-          MultiplyAdd(acc[VW2*w+10], xlm[kl], avec.sA);
-          MultiplyAdd(acc[VW2*w+11], xlm[kl], avec.sB);
-          MultiplyAdd(acc[VW2*w+12], xlm[kl], avec.sC);
-          MultiplyAdd(acc[VW2*w+13], xlm[kl], avec.sD);
-          MultiplyAdd(acc[VW2*w+14], xlm[kl], avec.sE);
-          MultiplyAdd(acc[VW2*w+15], xlm[kl], avec.sF);
+          MultiplyAdd(acc[VW2*_w+0], xlm[_kl], avec.s0);
+          MultiplyAdd(acc[VW2*_w+1], xlm[_kl], avec.s1);
+          MultiplyAdd(acc[VW2*_w+2], xlm[_kl], avec.s2);
+          MultiplyAdd(acc[VW2*_w+3], xlm[_kl], avec.s3);
+          MultiplyAdd(acc[VW2*_w+4], xlm[_kl], avec.s4);
+          MultiplyAdd(acc[VW2*_w+5], xlm[_kl], avec.s5);
+          MultiplyAdd(acc[VW2*_w+6], xlm[_kl], avec.s6);
+          MultiplyAdd(acc[VW2*_w+7], xlm[_kl], avec.s7);
+          MultiplyAdd(acc[VW2*_w+8], xlm[_kl], avec.s8);
+          MultiplyAdd(acc[VW2*_w+9], xlm[_kl], avec.s9);
+          MultiplyAdd(acc[VW2*_w+10], xlm[_kl], avec.sA);
+          MultiplyAdd(acc[VW2*_w+11], xlm[_kl], avec.sB);
+          MultiplyAdd(acc[VW2*_w+12], xlm[_kl], avec.sC);
+          MultiplyAdd(acc[VW2*_w+13], xlm[_kl], avec.sD);
+          MultiplyAdd(acc[VW2*_w+14], xlm[_kl], avec.sE);
+          MultiplyAdd(acc[VW2*_w+15], xlm[_kl], avec.sF);
         #endif
       }
     }
@@ -175,10 +176,10 @@ void XgemvFast(const int m, const int n,
 
   // Stores the final result
   #pragma unroll
-  for (int w=0; w<WPT2; ++w) {
-    const int gid = WPT2*get_global_id(0) + w;
+  for (int _w = 0; _w < WPT2; _w += 1) {
+    const int gid = WPT2*get_global_id(0) + _w;
     real yval = ygm[gid*y_inc + y_offset];
-    AXPBY(ygm[gid*y_inc + y_offset], alpha, acc[w], beta, yval);
+    AXPBY(ygm[gid*y_inc + y_offset], alpha, acc[_w], beta, yval);
   }
 }
 
@@ -226,46 +227,46 @@ void XgemvFastRot(const int m, const int n,
 
     // Loads the matrix A into local memory
     #pragma unroll
-    for (int kl=0; kl<WPT3/VW3; ++kl) {
+    for (int _kl = 0; _kl < WPT3/VW3; _kl += 1) {
       const int x = (kwg/VW3) + lid_mod;
-      const int y = get_group_id(0) * WGS3 + lid_div * (WPT3/VW3) + kl;
+      const int y = get_group_id(0) * WGS3 + lid_div * (WPT3/VW3) + _kl;
       realVFR avec = agm[(a_ld/VW3) * y + x];
       #if VW3 == 1
-        tile[kl*VW3 + 0][lid] = avec;
+        tile[_kl*VW3 + 0][lid] = avec;
       #elif VW3 == 2
-        tile[kl*VW3 + 0][lid] = avec.x;
-        tile[kl*VW3 + 1][lid] = avec.y;
+        tile[_kl*VW3 + 0][lid] = avec.x;
+        tile[_kl*VW3 + 1][lid] = avec.y;
       #elif VW3 == 4
-        tile[kl*VW3 + 0][lid] = avec.x;
-        tile[kl*VW3 + 1][lid] = avec.y;
-        tile[kl*VW3 + 2][lid] = avec.z;
-        tile[kl*VW3 + 3][lid] = avec.w;
+        tile[_kl*VW3 + 0][lid] = avec.x;
+        tile[_kl*VW3 + 1][lid] = avec.y;
+        tile[_kl*VW3 + 2][lid] = avec.z;
+        tile[_kl*VW3 + 3][lid] = avec.w;
       #elif VW3 == 8
-        tile[kl*VW3 + 0][lid] = avec.s0;
-        tile[kl*VW3 + 1][lid] = avec.s1;
-        tile[kl*VW3 + 2][lid] = avec.s2;
-        tile[kl*VW3 + 3][lid] = avec.s3;
-        tile[kl*VW3 + 4][lid] = avec.s4;
-        tile[kl*VW3 + 5][lid] = avec.s5;
-        tile[kl*VW3 + 6][lid] = avec.s6;
-        tile[kl*VW3 + 7][lid] = avec.s7;
+        tile[_kl*VW3 + 0][lid] = avec.s0;
+        tile[_kl*VW3 + 1][lid] = avec.s1;
+        tile[_kl*VW3 + 2][lid] = avec.s2;
+        tile[_kl*VW3 + 3][lid] = avec.s3;
+        tile[_kl*VW3 + 4][lid] = avec.s4;
+        tile[_kl*VW3 + 5][lid] = avec.s5;
+        tile[_kl*VW3 + 6][lid] = avec.s6;
+        tile[_kl*VW3 + 7][lid] = avec.s7;
       #elif VW3 == 16
-        tile[kl*VW3 + 0][lid] = avec.s0;
-        tile[kl*VW3 + 1][lid] = avec.s1;
-        tile[kl*VW3 + 2][lid] = avec.s2;
-        tile[kl*VW3 + 3][lid] = avec.s3;
-        tile[kl*VW3 + 4][lid] = avec.s4;
-        tile[kl*VW3 + 5][lid] = avec.s5;
-        tile[kl*VW3 + 6][lid] = avec.s6;
-        tile[kl*VW3 + 7][lid] = avec.s7;
-        tile[kl*VW3 + 8][lid] = avec.s8;
-        tile[kl*VW3 + 9][lid] = avec.s9;
-        tile[kl*VW3 + 10][lid] = avec.sA;
-        tile[kl*VW3 + 11][lid] = avec.sB;
-        tile[kl*VW3 + 12][lid] = avec.sC;
-        tile[kl*VW3 + 13][lid] = avec.sD;
-        tile[kl*VW3 + 14][lid] = avec.sE;
-        tile[kl*VW3 + 15][lid] = avec.sF;
+        tile[_kl*VW3 + 0][lid] = avec.s0;
+        tile[_kl*VW3 + 1][lid] = avec.s1;
+        tile[_kl*VW3 + 2][lid] = avec.s2;
+        tile[_kl*VW3 + 3][lid] = avec.s3;
+        tile[_kl*VW3 + 4][lid] = avec.s4;
+        tile[_kl*VW3 + 5][lid] = avec.s5;
+        tile[_kl*VW3 + 6][lid] = avec.s6;
+        tile[_kl*VW3 + 7][lid] = avec.s7;
+        tile[_kl*VW3 + 8][lid] = avec.s8;
+        tile[_kl*VW3 + 9][lid] = avec.s9;
+        tile[_kl*VW3 + 10][lid] = avec.sA;
+        tile[_kl*VW3 + 11][lid] = avec.sB;
+        tile[_kl*VW3 + 12][lid] = avec.sC;
+        tile[_kl*VW3 + 13][lid] = avec.sD;
+        tile[_kl*VW3 + 14][lid] = avec.sE;
+        tile[_kl*VW3 + 15][lid] = avec.sF;
       #endif
     }
 
@@ -274,11 +275,11 @@ void XgemvFastRot(const int m, const int n,
 
     // The multiply-add function (rotated)
     #pragma unroll
-    for (int kl=0; kl<WPT3/VW3; ++kl) {
+    for (int _kl = 0; _kl < WPT3/VW3; _kl += 1) {
       #pragma unroll
-      for (int v=0; v<VW3; ++v) {
-        real aval = tile[lid_mod*VW3 + v][lid_div * (WPT3/VW3) + kl];
-        real xval = xlm[kl*VW3 + v];
+      for (int _v = 0; _v < VW3; _v += 1) {
+        real aval = tile[lid_mod*VW3 + _v][lid_div * (WPT3/VW3) + _kl];
+        real xval = xlm[_kl*VW3 + _v];
         MultiplyAdd(acc, xval, aval);
       }
     }
diff --git a/src/kernels/level2/xger.opencl b/src/kernels/level2/xger.opencl
index 1b9ded12..ca6071cd 100644
--- a/src/kernels/level2/xger.opencl
+++ b/src/kernels/level2/xger.opencl
@@ -28,7 +28,9 @@ void Xger(const int max1, const int max2,
   const real alpha = GetRealArg(arg_alpha);
 
   // Register storage for X and Y
+  #pragma promote_to_registers
   real xvalues[WPT];
+  #pragma promote_to_registers
   real yvalues[WPT];
 
   // Row-major version
@@ -36,31 +38,31 @@ void Xger(const int max1, const int max2,
 
     // Loads the X-vector
     #pragma unroll
-    for (int w=0; w<WPT; ++w) {
-      const int id2 = w*get_global_size(1) + get_global_id(1);
-      xvalues[w] = LoadVector(id2, max2, xgm, x_offset, x_inc, false);
+    for (int _w = 0; _w < WPT; _w += 1) {
+      const int id2 = _w*get_global_size(1) + get_global_id(1);
+      xvalues[_w] = LoadVector(id2, max2, xgm, x_offset, x_inc, false);
     }
 
     // Loads the Y-vector
     #pragma unroll
-    for (int w=0; w<WPT; ++w) {
-      const int id1 = w*get_global_size(0) + get_global_id(0);
-      yvalues[w] = LoadVector(id1, max1, ygm, y_offset, y_inc, true);
+    for (int _w = 0; _w < WPT; _w += 1) {
+      const int id1 = _w*get_global_size(0) + get_global_id(0);
+      yvalues[_w] = LoadVector(id1, max1, ygm, y_offset, y_inc, true);
     }
 
     // Loops over the work per thread twice
     #pragma unroll
-    for (int w1=0; w1<WPT; ++w1) {
+    for (int _w1 = 0; _w1 < WPT; _w1 += 1) {
       #pragma unroll
-      for (int w2=0; w2<WPT; ++w2) {
+      for (int _w2 = 0; _w2 < WPT; _w2 += 1) {
 
         // Global thread IDs
-        const int id1 = w1*get_global_size(0) + get_global_id(0);
-        const int id2 = w2*get_global_size(1) + get_global_id(1);
+        const int id1 = _w1*get_global_size(0) + get_global_id(0);
+        const int id2 = _w2*get_global_size(1) + get_global_id(1);
 
         // Loads A, performs the operation, and stores the result into A
         MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld,
-                     alpha, xvalues[w2], yvalues[w1], false);
+                     alpha, xvalues[_w2], yvalues[_w1], false);
       }
     }
   }
@@ -70,31 +72,31 @@ void Xger(const int max1, const int max2,
 
     // Loads the X-vector
     #pragma unroll
-    for (int w=0; w<WPT; ++w) {
-      const int id1 = w*get_global_size(0) + get_global_id(0);
-      xvalues[w] = LoadVector(id1, max1, xgm, x_offset, x_inc, false);
+    for (int _w = 0; _w < WPT; _w += 1) {
+      const int id1 = _w*get_global_size(0) + get_global_id(0);
+      xvalues[_w] = LoadVector(id1, max1, xgm, x_offset, x_inc, false);
     }
 
     // Loads the Y-vector
     #pragma unroll
-    for (int w=0; w<WPT; ++w) {
-      const int id2 = w*get_global_size(1) + get_global_id(1);
-      yvalues[w] = LoadVector(id2, max2, ygm, y_offset, y_inc, true);
+    for (int _w = 0; _w < WPT; _w += 1) {
+      const int id2 = _w*get_global_size(1) + get_global_id(1);
+      yvalues[_w] = LoadVector(id2, max2, ygm, y_offset, y_inc, true);
     }
 
     // Loops over the work per thread twice
     #pragma unroll
-    for (int w1=0; w1<WPT; ++w1) {
+    for (int _w1 = 0; _w1 < WPT; _w1 += 1) {
       #pragma unroll
-      for (int w2=0; w2<WPT; ++w2) {
+      for (int _w2 = 0; _w2 < WPT; _w2 += 1) {
 
         // Global thread IDs
-        const int id1 = w1*get_global_size(0) + get_global_id(0);
-        const int id2 = w2*get_global_size(1) + get_global_id(1);
+        const int id1 = _w1*get_global_size(0) + get_global_id(0);
+        const int id2 = _w2*get_global_size(1) + get_global_id(1);
 
         // Loads A, performs the operation, and stores the result into A
         MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld,
-                     alpha, xvalues[w1], yvalues[w2], false);
+                     alpha, xvalues[_w1], yvalues[_w2], false);
       }
     }
   }
diff --git a/src/kernels/level3/copy_fast.opencl b/src/kernels/level3/copy_fast.opencl
index 695b9003..ef8a9017 100644
--- a/src/kernels/level3/copy_fast.opencl
+++ b/src/kernels/level3/copy_fast.opencl
@@ -42,9 +42,9 @@ void CopyMatrixFast(const int ld,
                     const real_arg arg_alpha) {
   const real alpha = GetRealArg(arg_alpha);
   #pragma unroll
-  for (int w_one=0; w_one<COPY_WPT; ++w_one) {
+  for (int _w_one = 0; _w_one < COPY_WPT; _w_one += 1) {
     const int id_one = get_global_id(0);
-    const int id_two = (get_group_id(1)*COPY_WPT + w_one) * COPY_DIMY + get_local_id(1);
+    const int id_two = (get_group_id(1)*COPY_WPT + _w_one) * COPY_DIMY + get_local_id(1);
     const int id = id_two*(ld/COPY_VW) + id_one;
     realC result;
     #if COPY_VW == 1
diff --git a/src/kernels/level3/copy_pad.opencl b/src/kernels/level3/copy_pad.opencl
index 6eeadbd1..2e191514 100644
--- a/src/kernels/level3/copy_pad.opencl
+++ b/src/kernels/level3/copy_pad.opencl
@@ -35,11 +35,11 @@ INLINE_FUNC void _CopyPadMatrix(const int src_one, const int src_two,
 
   // Loops over the work per thread in both dimensions
   #pragma unroll
-  for (int w_one=0; w_one<PAD_WPTX; ++w_one) {
-    const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0);
+  for (int _w_one = 0; _w_one < PAD_WPTX; _w_one += 1) {
+    const int id_one = (get_group_id(0)*PAD_WPTX + _w_one) * PAD_DIMX + get_local_id(0);
     #pragma unroll
-    for (int w_two=0; w_two<PAD_WPTY; ++w_two) {
-      const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1);
+    for (int _w_two = 0; _w_two < PAD_WPTY; _w_two += 1) {
+      const int id_two = (get_group_id(1)*PAD_WPTY + _w_two) * PAD_DIMY + get_local_id(1);
       if (id_two < dest_two && id_one < dest_one) {
 
         // Loads data if the thread IDs are within bounds of the source matrix. Otherwise, set the
@@ -91,11 +91,11 @@ INLINE_FUNC void _CopyMatrix(const int src_one, const int src_two,
 
   // Loops over the work per thread in both dimensions
   #pragma unroll
-  for (int w_one=0; w_one<PAD_WPTX; ++w_one) {
-    const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0);
+  for (int _w_one = 0; _w_one < PAD_WPTX; _w_one += 1) {
+    const int id_one = (get_group_id(0)*PAD_WPTX + _w_one) * PAD_DIMX + get_local_id(0);
     #pragma unroll
-    for (int w_two=0; w_two<PAD_WPTY; ++w_two) {
-      const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1);
+    for (int _w_two = 0; _w_two < PAD_WPTY; _w_two += 1) {
+      const int id_two = (get_group_id(1)*PAD_WPTY + _w_two) * PAD_DIMY + get_local_id(1);
 
       // Masking in case of triangular matrices: updates only the upper or lower part
       bool condition = true;