1 files changed, 22 insertions, 18 deletions

diff --git a/src/kernels/level2/xher2.opencl b/src/kernels/level2/xher2.opencl
index 00a756c9..73305149 100644
--- a/src/kernels/level2/xher2.opencl
+++ b/src/kernels/level2/xher2.opencl
@@ -28,37 +28,41 @@ void Xher2(const int n,
   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];
+  #pragma promote_to_registers
   real xtvalues[WPT];
+  #pragma promote_to_registers
   real ytvalues[WPT];
 
   // 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, n, xgm, x_offset, x_inc, !is_rowmajor);
+  for (int _w = 0; _w < WPT; _w += 1) {
+    const int id2 = _w*get_global_size(1) + get_global_id(1);
+    xvalues[_w] = LoadVector(id2, n, xgm, x_offset, x_inc, !is_rowmajor);
   }
 
   // Loads the X-transposed-vector
   #pragma unroll
-  for (int w=0; w<WPT; ++w) {
-    const int id1 = w*get_global_size(0) + get_global_id(0);
-    xtvalues[w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor);
+  for (int _w = 0; _w < WPT; _w += 1) {
+    const int id1 = _w*get_global_size(0) + get_global_id(0);
+    xtvalues[_w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor);
   }
 
   // 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, n, ygm, y_offset, y_inc, is_rowmajor);
+  for (int _w = 0; _w < WPT; _w += 1) {
+    const int id1 = _w*get_global_size(0) + get_global_id(0);
+    yvalues[_w] = LoadVector(id1, n, ygm, y_offset, y_inc, is_rowmajor);
   }
 
   // Loads the Y-transposed-vector
   #pragma unroll
-  for (int w=0; w<WPT; ++w) {
-    const int id2 = w*get_global_size(1) + get_global_id(1);
-    ytvalues[w] = LoadVector(id2, n, ygm, y_offset, y_inc, !is_rowmajor);
+  for (int _w = 0; _w < WPT; _w += 1) {
+    const int id2 = _w*get_global_size(1) + get_global_id(1);
+    ytvalues[_w] = LoadVector(id2, n, ygm, y_offset, y_inc, !is_rowmajor);
   }
 
   // Sets the proper value of alpha in case conjugation is needed
@@ -75,13 +79,13 @@ void Xher2(const int n,
 
   // 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);
 
       // Skip these threads if they do not contain threads contributing to the matrix-triangle
       if ((is_upper && (id1 > id2)) || (!is_upper && (id2 > id1))) {
@@ -91,8 +95,8 @@ void Xher2(const int n,
       // Loads A, performs the operation, and stores the result into A
       else {
         MatrixUpdate2(id1, id2, n, n, agm, a_offset, a_ld,
-                      alpha1, xvalues[w2], yvalues[w1],
-                      alpha2, xtvalues[w1], ytvalues[w2], is_upper);
+                      alpha1, xvalues[_w2], yvalues[_w1],
+                      alpha2, xtvalues[_w1], ytvalues[_w2], is_upper);
       }
     }
   }