Improved array-to-register promotion, now handling function calls as well

2 files changed, 45 insertions, 45 deletions

diff --git a/src/kernels/level2/xgemv.opencl b/src/kernels/level2/xgemv.opencl
index 2a50e8fb..ba29aba6 100644
--- a/src/kernels/level2/xgemv.opencl
+++ b/src/kernels/level2/xgemv.opencl
@@ -228,10 +228,10 @@ void Xgemv(const int m, const int n,
 
   // Initializes the accumulation register
   #pragma promote_to_registers
-  real acc[WPT1];
+  real acc1[WPT1];
   #pragma unroll
   for (int _w = 0; _w < WPT1; _w += 1) {
-    SetToZero(acc[_w]);
+    SetToZero(acc1[_w]);
   }
 
   // Divides the work in a main and tail section
@@ -262,7 +262,7 @@ void Xgemv(const int m, const int n,
               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(acc1[_w], xlm[kloop + _kunroll], value);
             }
           }
         }
@@ -273,7 +273,7 @@ void Xgemv(const int m, const int n,
               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(acc1[_w], xlm[kloop + _kunroll], value);
             }
           }
         }
@@ -295,20 +295,20 @@ void Xgemv(const int m, const int n,
         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(acc1[_w], xgm[k*x_inc + x_offset], value);
         }
       }
       else { // Transposed
         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(acc1[_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, acc1[_w], beta, yval);
     }
   }
 }
diff --git a/src/kernels/level2/xgemv_fast.opencl b/src/kernels/level2/xgemv_fast.opencl
index 892bc55c..45ceb36c 100644
--- a/src/kernels/level2/xgemv_fast.opencl
+++ b/src/kernels/level2/xgemv_fast.opencl
@@ -106,10 +106,10 @@ void XgemvFast(const int m, const int n,
 
   // Initializes the accumulation registers
   #pragma promote_to_registers
-  real acc[WPT2];
+  real acc2[WPT2];
   #pragma unroll
   for (int _w = 0; _w < WPT2; _w += 1) {
-    SetToZero(acc[_w]);
+    SetToZero(acc2[_w]);
   }
 
   // Loops over work-group sized portions of the work
@@ -131,41 +131,41 @@ void XgemvFast(const int m, const int n,
         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(acc2[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(acc2[VW2*_w+0], xlm[_kl], avec.x);
+          MultiplyAdd(acc2[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(acc2[VW2*_w+0], xlm[_kl], avec.x);
+          MultiplyAdd(acc2[VW2*_w+1], xlm[_kl], avec.y);
+          MultiplyAdd(acc2[VW2*_w+2], xlm[_kl], avec.z);
+          MultiplyAdd(acc2[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(acc2[VW2*_w+0], xlm[_kl], avec.s0);
+          MultiplyAdd(acc2[VW2*_w+1], xlm[_kl], avec.s1);
+          MultiplyAdd(acc2[VW2*_w+2], xlm[_kl], avec.s2);
+          MultiplyAdd(acc2[VW2*_w+3], xlm[_kl], avec.s3);
+          MultiplyAdd(acc2[VW2*_w+4], xlm[_kl], avec.s4);
+          MultiplyAdd(acc2[VW2*_w+5], xlm[_kl], avec.s5);
+          MultiplyAdd(acc2[VW2*_w+6], xlm[_kl], avec.s6);
+          MultiplyAdd(acc2[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(acc2[VW2*_w+0], xlm[_kl], avec.s0);
+          MultiplyAdd(acc2[VW2*_w+1], xlm[_kl], avec.s1);
+          MultiplyAdd(acc2[VW2*_w+2], xlm[_kl], avec.s2);
+          MultiplyAdd(acc2[VW2*_w+3], xlm[_kl], avec.s3);
+          MultiplyAdd(acc2[VW2*_w+4], xlm[_kl], avec.s4);
+          MultiplyAdd(acc2[VW2*_w+5], xlm[_kl], avec.s5);
+          MultiplyAdd(acc2[VW2*_w+6], xlm[_kl], avec.s6);
+          MultiplyAdd(acc2[VW2*_w+7], xlm[_kl], avec.s7);
+          MultiplyAdd(acc2[VW2*_w+8], xlm[_kl], avec.s8);
+          MultiplyAdd(acc2[VW2*_w+9], xlm[_kl], avec.s9);
+          MultiplyAdd(acc2[VW2*_w+10], xlm[_kl], avec.sA);
+          MultiplyAdd(acc2[VW2*_w+11], xlm[_kl], avec.sB);
+          MultiplyAdd(acc2[VW2*_w+12], xlm[_kl], avec.sC);
+          MultiplyAdd(acc2[VW2*_w+13], xlm[_kl], avec.sD);
+          MultiplyAdd(acc2[VW2*_w+14], xlm[_kl], avec.sE);
+          MultiplyAdd(acc2[VW2*_w+15], xlm[_kl], avec.sF);
         #endif
       }
     }
@@ -179,7 +179,7 @@ void XgemvFast(const int m, const int n,
   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, acc2[_w], beta, yval);
   }
 }
 
@@ -214,8 +214,8 @@ void XgemvFastRot(const int m, const int n,
   __local real xlm[WPT3];
 
   // Initializes the accumulation register
-  real acc;
-  SetToZero(acc);
+  real acc3;
+  SetToZero(acc3);
 
   // Loops over tile-sized portions of the work
   for (int kwg=0; kwg<n; kwg+=WPT3) {
@@ -280,7 +280,7 @@ void XgemvFastRot(const int m, const int n,
       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);
+        MultiplyAdd(acc3, xval, aval);
       }
     }
 
@@ -291,7 +291,7 @@ void XgemvFastRot(const int m, const int n,
   // Stores the final result
   const int gid = get_global_id(0);
   real yval = ygm[gid * y_inc + y_offset];
-  AXPBY(ygm[gid * y_inc + y_offset], alpha, acc, beta, yval);
+  AXPBY(ygm[gid * y_inc + y_offset], alpha, acc3, beta, yval);
 }
 
 // =================================================================================================