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_pad.opencl | 48 ++++++++++++++++-----------------
 1 file changed, 24 insertions(+), 24 deletions(-)

(limited to 'src/kernels/level3/transpose_pad.opencl')

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