diff options
Diffstat (limited to 'src/kernels/level3')
-rw-r--r-- | src/kernels/level3/transpose_fast.opencl | 78 | ||||
-rw-r--r-- | src/kernels/level3/transpose_pad.opencl | 48 |
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); |