summaryrefslogtreecommitdiff
path: root/src/kernels/level3
diff options
context:
space:
mode:
authorCedric Nugteren <web@cedricnugteren.nl>2017-12-03 12:00:37 +0100
committerCedric Nugteren <web@cedricnugteren.nl>2017-12-03 12:00:37 +0100
commit60312e5878fd45225158dd8545a01366f937a871 (patch)
tree0755588371b125df6f5c3a9e4be04a0ca53bdd63 /src/kernels/level3
parent92842024b0e9fb2df8e8e43c4499d0d2d13fefc0 (diff)
Reformated transpose kernels for the pre-processor; extended the amount of tests
Diffstat (limited to 'src/kernels/level3')
-rw-r--r--src/kernels/level3/transpose_fast.opencl78
-rw-r--r--src/kernels/level3/transpose_pad.opencl48
2 files changed, 64 insertions, 62 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);