From 60312e5878fd45225158dd8545a01366f937a871 Mon Sep 17 00:00:00 2001
From: Cedric Nugteren <web@cedricnugteren.nl>
Date: Sun, 3 Dec 2017 12:00:37 +0100
Subject: Reformated transpose kernels for the pre-processor; extended the
 amount of tests

---
 src/kernels/level3/transpose_fast.opencl | 78 ++++++++++++++++----------------
 src/kernels/level3/transpose_pad.opencl  | 48 ++++++++++----------
 test/correctness/misc/preprocessor.cpp   | 25 ++++++++--
 3 files changed, 85 insertions(+), 66 deletions(-)

diff --git a/src/kernels/level3/transpose_fast.opencl b/src/kernels/level3/transpose_fast.opencl
index 37b25d99..8fa7405c 100644
--- a/src/kernels/level3/transpose_fast.opencl
+++ b/src/kernels/level3/transpose_fast.opencl
@@ -57,29 +57,31 @@ void TransposeMatrixFast(const int ld,
 
   // Loops over the work per thread
   #pragma unroll
-  for (int w_one=0; w_one<TRA_WPT; ++w_one) {
+  for (int _w_one = 0; _w_one < TRA_WPT; _w_one += 1) {
 
     // Computes the identifiers for the source matrix. Note that the local and global dimensions
     // do not correspond to each other!
     const int id_one = gid1 * TRA_DIM + get_local_id(0);
-    const int id_two = (gid0 * TRA_DIM + get_local_id(1))*TRA_WPT + w_one;
+    const int id_two = (gid0 * TRA_DIM + get_local_id(1))*TRA_WPT + _w_one;
 
     // Loads data into the local memory
     realT value = src[id_two*(ld/TRA_WPT) + id_one];
-    tile[get_local_id(0)*TRA_WPT + w_one][get_local_id(1)] = value;
+    tile[get_local_id(0)*TRA_WPT + _w_one][get_local_id(1)] = value;
   }
 
   // Synchronizes all threads in a workgroup
   barrier(CLK_LOCAL_MEM_FENCE);
 
   // Loads transposed data from the local memory
+  #pragma promote_to_registers
   realT v[TRA_WPT];
   #pragma unroll
-  for (int w_one=0; w_one<TRA_WPT; ++w_one) {
-    v[w_one] = tile[get_local_id(1)*TRA_WPT + w_one][get_local_id(0)];
+  for (int _w_one = 0; _w_one < TRA_WPT; _w_one += 1) {
+    v[_w_one] = tile[get_local_id(1)*TRA_WPT + _w_one][get_local_id(0)];
   }
 
   // Performs the register-level transpose of the vectorized data
+  #pragma promote_to_registers
   realT results[TRA_WPT];
   #if TRA_WPT == 1
     results[0] = v[0];
@@ -121,47 +123,47 @@ void TransposeMatrixFast(const int ld,
 
   // Multiplies by alpha and then stores the results into the destination matrix
   #pragma unroll
-  for (int w_two=0; w_two<TRA_WPT; ++w_two) {
+  for (int _w_two = 0; _w_two < TRA_WPT; _w_two += 1) {
     realT result;
     #if TRA_WPT == 1
-      Multiply(result, alpha, results[w_two]);
+      Multiply(result, alpha, results[_w_two]);
     #elif TRA_WPT == 2
-      Multiply(result.x, alpha, results[w_two].x);
-      Multiply(result.y, alpha, results[w_two].y);
+      Multiply(result.x, alpha, results[_w_two].x);
+      Multiply(result.y, alpha, results[_w_two].y);
     #elif TRA_WPT == 4
-      Multiply(result.x, alpha, results[w_two].x);
-      Multiply(result.y, alpha, results[w_two].y);
-      Multiply(result.z, alpha, results[w_two].z);
-      Multiply(result.w, alpha, results[w_two].w);
+      Multiply(result.x, alpha, results[_w_two].x);
+      Multiply(result.y, alpha, results[_w_two].y);
+      Multiply(result.z, alpha, results[_w_two].z);
+      Multiply(result.w, alpha, results[_w_two].w);
     #elif TRA_WPT == 8
-      Multiply(result.s0, alpha, results[w_two].s0);
-      Multiply(result.s1, alpha, results[w_two].s1);
-      Multiply(result.s2, alpha, results[w_two].s2);
-      Multiply(result.s3, alpha, results[w_two].s3);
-      Multiply(result.s4, alpha, results[w_two].s4);
-      Multiply(result.s5, alpha, results[w_two].s5);
-      Multiply(result.s6, alpha, results[w_two].s6);
-      Multiply(result.s7, alpha, results[w_two].s7);
+      Multiply(result.s0, alpha, results[_w_two].s0);
+      Multiply(result.s1, alpha, results[_w_two].s1);
+      Multiply(result.s2, alpha, results[_w_two].s2);
+      Multiply(result.s3, alpha, results[_w_two].s3);
+      Multiply(result.s4, alpha, results[_w_two].s4);
+      Multiply(result.s5, alpha, results[_w_two].s5);
+      Multiply(result.s6, alpha, results[_w_two].s6);
+      Multiply(result.s7, alpha, results[_w_two].s7);
     #elif TRA_WPT == 16
-      Multiply(result.s0, alpha, results[w_two].s0);
-      Multiply(result.s1, alpha, results[w_two].s1);
-      Multiply(result.s2, alpha, results[w_two].s2);
-      Multiply(result.s3, alpha, results[w_two].s3);
-      Multiply(result.s4, alpha, results[w_two].s4);
-      Multiply(result.s5, alpha, results[w_two].s5);
-      Multiply(result.s6, alpha, results[w_two].s6);
-      Multiply(result.s7, alpha, results[w_two].s7);
-      Multiply(result.s8, alpha, results[w_two].s8);
-      Multiply(result.s9, alpha, results[w_two].s9);
-      Multiply(result.sA, alpha, results[w_two].sA);
-      Multiply(result.sB, alpha, results[w_two].sB);
-      Multiply(result.sC, alpha, results[w_two].sC);
-      Multiply(result.sD, alpha, results[w_two].sD);
-      Multiply(result.sE, alpha, results[w_two].sE);
-      Multiply(result.sF, alpha, results[w_two].sF);
+      Multiply(result.s0, alpha, results[_w_two].s0);
+      Multiply(result.s1, alpha, results[_w_two].s1);
+      Multiply(result.s2, alpha, results[_w_two].s2);
+      Multiply(result.s3, alpha, results[_w_two].s3);
+      Multiply(result.s4, alpha, results[_w_two].s4);
+      Multiply(result.s5, alpha, results[_w_two].s5);
+      Multiply(result.s6, alpha, results[_w_two].s6);
+      Multiply(result.s7, alpha, results[_w_two].s7);
+      Multiply(result.s8, alpha, results[_w_two].s8);
+      Multiply(result.s9, alpha, results[_w_two].s9);
+      Multiply(result.sA, alpha, results[_w_two].sA);
+      Multiply(result.sB, alpha, results[_w_two].sB);
+      Multiply(result.sC, alpha, results[_w_two].sC);
+      Multiply(result.sD, alpha, results[_w_two].sD);
+      Multiply(result.sE, alpha, results[_w_two].sE);
+      Multiply(result.sF, alpha, results[_w_two].sF);
     #endif
     const int id_one = gid0*TRA_DIM + get_local_id(0);
-    const int id_two = (gid1*TRA_DIM + get_local_id(1))*TRA_WPT + w_two;
+    const int id_two = (gid1*TRA_DIM + get_local_id(1))*TRA_WPT + _w_two;
     dest[id_two*(ld/TRA_WPT) + id_one] = result;
   }
 }
diff --git a/src/kernels/level3/transpose_pad.opencl b/src/kernels/level3/transpose_pad.opencl
index ba9a6a56..67c2bf72 100644
--- a/src/kernels/level3/transpose_pad.opencl
+++ b/src/kernels/level3/transpose_pad.opencl
@@ -36,14 +36,14 @@ INLINE_FUNC void _TransposePadMatrix(LOCAL_PTR real* tile,
 
   // Loop over the work per thread
   #pragma unroll
-  for (int w_one=0; w_one<PADTRA_WPT; ++w_one) {
+  for (int _w_one = 0; _w_one < PADTRA_WPT; _w_one += 1) {
     #pragma unroll
-    for (int w_two=0; w_two<PADTRA_WPT; ++w_two) {
+    for (int _w_two = 0; _w_two < PADTRA_WPT; _w_two += 1) {
 
       // Computes the identifiers for the source matrix. Note that the local and global dimensions
       // do not correspond to each other!
-      const int id_src_one = (get_group_id(1)*PADTRA_WPT + w_two) * PADTRA_TILE + get_local_id(0);
-      const int id_src_two = (get_group_id(0)*PADTRA_WPT + w_one) * PADTRA_TILE + get_local_id(1);
+      const int id_src_one = (get_group_id(1)*PADTRA_WPT + _w_two) * PADTRA_TILE + get_local_id(0);
+      const int id_src_two = (get_group_id(0)*PADTRA_WPT + _w_one) * PADTRA_TILE + get_local_id(1);
 
       // Loads data into the local memory if the thread IDs are within bounds of the source matrix.
       // Otherwise, set the local memory value to zero.
@@ -52,8 +52,8 @@ INLINE_FUNC void _TransposePadMatrix(LOCAL_PTR real* tile,
       if (id_src_two < src_two && id_src_one < src_one) {
         value = src[id_src_two*src_ld + id_src_one + src_offset];
       }
-      const int tile_id0 = get_local_id(0)*PADTRA_WPT + w_one;
-      const int tile_id1 = get_local_id(1)*PADTRA_WPT + w_two;
+      const int tile_id0 = get_local_id(0)*PADTRA_WPT + _w_one;
+      const int tile_id1 = get_local_id(1)*PADTRA_WPT + _w_two;
       tile[tile_id1 * (PADTRA_WPT*PADTRA_TILE + PADTRA_PAD) + tile_id0] = value;
     }
   }
@@ -63,18 +63,18 @@ INLINE_FUNC void _TransposePadMatrix(LOCAL_PTR real* tile,
 
   // Loop over the work per thread
   #pragma unroll
-  for (int w_one=0; w_one<PADTRA_WPT; ++w_one) {
+  for (int _w_one = 0; _w_one < PADTRA_WPT; _w_one += 1) {
     #pragma unroll
-    for (int w_two=0; w_two<PADTRA_WPT; ++w_two) {
+    for (int _w_two = 0; _w_two < PADTRA_WPT; _w_two += 1) {
 
       // Computes the identifiers for the destination matrix
-      const int id_dest_one = (get_group_id(0)*PADTRA_WPT + w_one) * PADTRA_TILE + get_local_id(0);
-      const int id_dest_two = (get_group_id(1)*PADTRA_WPT + w_two) * PADTRA_TILE + get_local_id(1);
+      const int id_dest_one = (get_group_id(0)*PADTRA_WPT + _w_one) * PADTRA_TILE + get_local_id(0);
+      const int id_dest_two = (get_group_id(1)*PADTRA_WPT + _w_two) * PADTRA_TILE + get_local_id(1);
 
       // Stores the transposed value in the destination matrix
       if ((id_dest_one < dest_one) && (id_dest_two < dest_two)) {
-        const int tile_id0 = get_local_id(1)*PADTRA_WPT + w_one;
-        const int tile_id1 = get_local_id(0)*PADTRA_WPT + w_two;
+        const int tile_id0 = get_local_id(1)*PADTRA_WPT + _w_one;
+        const int tile_id1 = get_local_id(0)*PADTRA_WPT + _w_two;
         real value = tile[tile_id1 * (PADTRA_WPT*PADTRA_TILE + PADTRA_PAD) + tile_id0];
         if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
         Multiply(dest[id_dest_two*dest_ld + id_dest_one + dest_offset], alpha, value);
@@ -118,20 +118,20 @@ INLINE_FUNC void _TransposeMatrix(LOCAL_PTR real* tile,
 
   // Loop over the work per thread
   #pragma unroll
-  for (int w_one=0; w_one<PADTRA_WPT; ++w_one) {
+  for (int _w_one = 0; _w_one < PADTRA_WPT; _w_one += 1) {
     #pragma unroll
-    for (int w_two=0; w_two<PADTRA_WPT; ++w_two) {
+    for (int _w_two = 0; _w_two < PADTRA_WPT; _w_two += 1) {
 
       // Computes the identifiers for the source matrix. Note that the local and global dimensions
       // do not correspond to each other!
-      const int id_src_one = (get_group_id(1)*PADTRA_WPT + w_two) * PADTRA_TILE + get_local_id(0);
-      const int id_src_two = (get_group_id(0)*PADTRA_WPT + w_one) * PADTRA_TILE + get_local_id(1);
+      const int id_src_one = (get_group_id(1)*PADTRA_WPT + _w_two) * PADTRA_TILE + get_local_id(0);
+      const int id_src_two = (get_group_id(0)*PADTRA_WPT + _w_one) * PADTRA_TILE + get_local_id(1);
 
       // Loads data into the local memory if the thread IDs are within bounds of the source matrix.
       if ((id_src_one < src_one) && (id_src_two < src_two)) {
         real value = src[id_src_two*src_ld + id_src_one + src_offset];
-        const int tile_id0 = get_local_id(0)*PADTRA_WPT + w_one;
-        const int tile_id1 = get_local_id(1)*PADTRA_WPT + w_two;
+        const int tile_id0 = get_local_id(0)*PADTRA_WPT + _w_one;
+        const int tile_id1 = get_local_id(1)*PADTRA_WPT + _w_two;
         tile[tile_id1 * (PADTRA_WPT*PADTRA_TILE + PADTRA_PAD) + tile_id0] = value;
       }
     }
@@ -142,13 +142,13 @@ INLINE_FUNC void _TransposeMatrix(LOCAL_PTR real* tile,
 
   // Loop over the work per thread
   #pragma unroll
-  for (int w_one=0; w_one<PADTRA_WPT; ++w_one) {
+  for (int _w_one = 0; _w_one < PADTRA_WPT; _w_one += 1) {
     #pragma unroll
-    for (int w_two=0; w_two<PADTRA_WPT; ++w_two) {
+    for (int _w_two = 0; _w_two < PADTRA_WPT; _w_two += 1) {
 
       // Computes the identifiers for the destination matrix
-      const int id_dest_one = (get_group_id(0)*PADTRA_WPT + w_one) * PADTRA_TILE + get_local_id(0);
-      const int id_dest_two = (get_group_id(1)*PADTRA_WPT + w_two) * PADTRA_TILE + get_local_id(1);
+      const int id_dest_one = (get_group_id(0)*PADTRA_WPT + _w_one) * PADTRA_TILE + get_local_id(0);
+      const int id_dest_two = (get_group_id(1)*PADTRA_WPT + _w_two) * PADTRA_TILE + get_local_id(1);
 
       // Masking in case of triangular matrices: updates only the upper or lower part
       bool condition = true;
@@ -160,8 +160,8 @@ INLINE_FUNC void _TransposeMatrix(LOCAL_PTR real* tile,
 
         // Stores the transposed value in the destination matrix
         if ((id_dest_one < dest_one) && (id_dest_two < dest_two)) {
-          const int tile_id0 = get_local_id(1)*PADTRA_WPT + w_one;
-          const int tile_id1 = get_local_id(0)*PADTRA_WPT + w_two;
+          const int tile_id0 = get_local_id(1)*PADTRA_WPT + _w_one;
+          const int tile_id1 = get_local_id(0)*PADTRA_WPT + _w_two;
           real value = tile[tile_id1 * (PADTRA_WPT*PADTRA_TILE + PADTRA_PAD) + tile_id0];
           if (diagonal_imag_zero == 1 && id_dest_one == id_dest_two) { ImagToZero(value); }
           Multiply(dest[id_dest_two*dest_ld + id_dest_one + dest_offset], alpha, value);
diff --git a/test/correctness/misc/preprocessor.cpp b/test/correctness/misc/preprocessor.cpp
index cea91116..d352ce37 100644
--- a/test/correctness/misc/preprocessor.cpp
+++ b/test/correctness/misc/preprocessor.cpp
@@ -133,6 +133,7 @@ size_t RunPreprocessor(int argc, char *argv[], const bool silent, const Precisio
     "#define WPT2 2\n"
     "#define WPT3 2\n"
     "#define UNROLL1 4\n"
+    "#define VW2 2\n"
     #include "../src/kernels/level2/xgemv.opencl"
     #include "../src/kernels/level2/xgemv_fast.opencl"
   ;
@@ -148,13 +149,29 @@ size_t RunPreprocessor(int argc, char *argv[], const bool silent, const Precisio
 
   // CopyPad
   const auto copy_pad_sources =
-      "#define PAD_WPTX 2\n"
-      "#define PAD_WPTY 2\n"
-#include "../src/kernels/level3/level3.opencl"
-#include "../src/kernels/level3/copy_pad.opencl"
+    "#define PAD_WPTX 2\n"
+    "#define PAD_WPTY 2\n"
+    #include "../src/kernels/level3/level3.opencl"
+    #include "../src/kernels/level3/copy_pad.opencl"
   ;
   if (TestKernel(device, context, "CopyPadMatrix", copy_pad_sources, precision)) { passed++; } else { errors++; }
 
+  // TransposeFast
+  const auto transpose_fast_sources =
+    "#define TRA_WPT 2\n"
+    #include "../src/kernels/level3/level3.opencl"
+    #include "../src/kernels/level3/transpose_fast.opencl"
+  ;
+  if (TestKernel(device, context, "TransposeMatrixFast", transpose_fast_sources, precision)) { passed++; } else { errors++; }
+
+  // TransposePad
+  const auto transpose_pad_sources =
+    "#define PADTRA_WPT 2\n"
+    #include "../src/kernels/level3/level3.opencl"
+    #include "../src/kernels/level3/transpose_pad.opencl"
+  ;
+  if (TestKernel(device, context, "TransposePadMatrix", transpose_pad_sources, precision)) { passed++; } else { errors++; }
+
 
 
   // Prints and returns the statistics
-- 
cgit v1.2.3