From 93ffb876c60838bee75d3bb25ebbcbfce02e2cc7 Mon Sep 17 00:00:00 2001
From: Cedric Nugteren <web@cedricnugteren.nl>
Date: Wed, 29 Nov 2017 20:21:08 +0100
Subject: Reformatted unrollable kernel loops and added the new
 promote_to_registers pragma for several kernels

---
 src/kernels/level2/xger.opencl | 46 ++++++++++++++++++++++--------------------
 1 file changed, 24 insertions(+), 22 deletions(-)

(limited to 'src/kernels/level2/xger.opencl')

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);
       }
     }
   }
-- 
cgit v1.2.3