diff options
-rw-r--r-- | src/kernel_preprocessor.cpp | 3 | ||||
-rw-r--r-- | src/kernels/level1/xaxpy.opencl | 8 | ||||
-rw-r--r-- | src/kernels/level2/xgemv.opencl | 34 | ||||
-rw-r--r-- | src/kernels/level2/xgemv_fast.opencl | 155 | ||||
-rw-r--r-- | src/kernels/level2/xger.opencl | 46 | ||||
-rw-r--r-- | src/kernels/level3/copy_fast.opencl | 4 | ||||
-rw-r--r-- | src/kernels/level3/copy_pad.opencl | 16 |
7 files changed, 133 insertions, 133 deletions
diff --git a/src/kernel_preprocessor.cpp b/src/kernel_preprocessor.cpp index b6c0b398..6aa853fd 100644 --- a/src/kernel_preprocessor.cpp +++ b/src/kernel_preprocessor.cpp @@ -164,9 +164,6 @@ std::vector<std::string> PreprocessDefinesAndComments(const std::string& source, if (HasOnlyDigits(value)) { defines.emplace(name, std::stoi(value)); } - else { - printf("'%s'\n", value.c_str()); - } } // Detect #ifndef blocks diff --git a/src/kernels/level1/xaxpy.opencl b/src/kernels/level1/xaxpy.opencl index 3a574ec2..74e49930 100644 --- a/src/kernels/level1/xaxpy.opencl +++ b/src/kernels/level1/xaxpy.opencl @@ -45,8 +45,8 @@ void XaxpyFaster(const int n, const real_arg arg_alpha, if (get_global_id(0) < n / (VW)) { #pragma unroll - for (int w = 0; w < WPT; w += 1) { - const int id = w*get_global_size(0) + get_global_id(0); + for (int _w = 0; _w < WPT; _w += 1) { + const int id = _w*get_global_size(0) + get_global_id(0); realV xvalue = xgm[id]; realV yvalue = ygm[id]; ygm[id] = MultiplyAddVector(yvalue, alpha, xvalue); @@ -63,8 +63,8 @@ void XaxpyFastest(const int n, const real_arg arg_alpha, const real alpha = GetRealArg(arg_alpha); #pragma unroll - for (int w = 0; w < WPT; w += 1) { - const int id = w*get_global_size(0) + get_global_id(0); + for (int _w = 0; _w < WPT; _w += 1) { + const int id = _w*get_global_size(0) + get_global_id(0); realV xvalue = xgm[id]; realV yvalue = ygm[id]; ygm[id] = MultiplyAddVector(yvalue, alpha, xvalue); diff --git a/src/kernels/level2/xgemv.opencl b/src/kernels/level2/xgemv.opencl index ea0478f0..2a50e8fb 100644 --- a/src/kernels/level2/xgemv.opencl +++ b/src/kernels/level2/xgemv.opencl @@ -227,10 +227,11 @@ void Xgemv(const int m, const int n, __local real xlm[WGS1]; // Initializes the accumulation register + #pragma promote_to_registers real acc[WPT1]; #pragma unroll - for (int w=0; w<WPT1; ++w) { - SetToZero(acc[w]); + for (int _w = 0; _w < WPT1; _w += 1) { + SetToZero(acc[_w]); } // Divides the work in a main and tail section @@ -248,30 +249,31 @@ void Xgemv(const int m, const int n, barrier(CLK_LOCAL_MEM_FENCE); // Loops over the work per thread, and checks whether in bounds - for (int w=0; w<WPT1; ++w) { - const int gid = w*get_global_size(0) + get_global_id(0); + #pragma unroll + for (int _w = 0; _w < WPT1; _w += 1) { + const int gid = _w*get_global_size(0) + get_global_id(0); if (gid < m) { // The multiply-add function for the main part (divisable by WGS1) if (a_rotated == 0) { // Not rotated for (int kloop=0; kloop<WGS1; kloop+=UNROLL1) { #pragma unroll - for (int kunroll=0; kunroll<UNROLL1; ++kunroll) { - const int k = kwg + kloop + kunroll; + for (int _kunroll = 0; _kunroll < UNROLL1; _kunroll += 1) { + 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(acc[_w], xlm[kloop + _kunroll], value); } } } else { // Transposed for (int kloop=0; kloop<WGS1; kloop+=UNROLL1) { #pragma unroll - for (int kunroll=0; kunroll<UNROLL1; ++kunroll) { - const int k = kwg + kloop + kunroll; + for (int _kunroll = 0; _kunroll < UNROLL1; _kunroll += 1) { + 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(acc[_w], xlm[kloop + _kunroll], value); } } } @@ -284,31 +286,29 @@ void Xgemv(const int m, const int n, // Loops over the work per thread, and checks whether in bounds #pragma unroll - for (int w=0; w<WPT1; ++w) { - const int gid = w*get_global_size(0) + get_global_id(0); + for (int _w = 0; _w < WPT1; _w += 1) { + const int gid = _w*get_global_size(0) + get_global_id(0); if (gid < m) { // The multiply-add function for the remainder part (not divisable by WGS1) if (a_rotated == 0) { // Not rotated - #pragma unroll 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(acc[_w], xgm[k*x_inc + x_offset], value); } } else { // Transposed - #pragma unroll 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(acc[_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, acc[_w], beta, yval); } } } diff --git a/src/kernels/level2/xgemv_fast.opencl b/src/kernels/level2/xgemv_fast.opencl index 8a08f076..892bc55c 100644 --- a/src/kernels/level2/xgemv_fast.opencl +++ b/src/kernels/level2/xgemv_fast.opencl @@ -105,10 +105,11 @@ void XgemvFast(const int m, const int n, __local real xlm[WGS2]; // Initializes the accumulation registers + #pragma promote_to_registers real acc[WPT2]; #pragma unroll - for (int w=0; w<WPT2; ++w) { - SetToZero(acc[w]); + for (int _w = 0; _w < WPT2; _w += 1) { + SetToZero(acc[_w]); } // Loops over work-group sized portions of the work @@ -123,48 +124,48 @@ void XgemvFast(const int m, const int n, // The multiply-add function (not rotated) #pragma unroll - for (int kl=0; kl<WGS2; ++kl) { - const int k = kwg + kl; + for (int _kl = 0; _kl < WGS2; _kl += 1) { + const int k = kwg + _kl; #pragma unroll - for (int w=0; w<WPT2/VW2; ++w) { - const int gid = (WPT2/VW2)*get_global_id(0) + w; + for (int _w = 0; _w < WPT2/VW2; _w += 1) { + 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(acc[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(acc[VW2*_w+0], xlm[_kl], avec.x); + MultiplyAdd(acc[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(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); #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(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); #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(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); #endif } } @@ -175,10 +176,10 @@ void XgemvFast(const int m, const int n, // Stores the final result #pragma unroll - for (int w=0; w<WPT2; ++w) { - const int gid = WPT2*get_global_id(0) + w; + 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, acc[_w], beta, yval); } } @@ -226,46 +227,46 @@ void XgemvFastRot(const int m, const int n, // Loads the matrix A into local memory #pragma unroll - for (int kl=0; kl<WPT3/VW3; ++kl) { + for (int _kl = 0; _kl < WPT3/VW3; _kl += 1) { const int x = (kwg/VW3) + lid_mod; - const int y = get_group_id(0) * WGS3 + lid_div * (WPT3/VW3) + kl; + const int y = get_group_id(0) * WGS3 + lid_div * (WPT3/VW3) + _kl; realVFR avec = agm[(a_ld/VW3) * y + x]; #if VW3 == 1 - tile[kl*VW3 + 0][lid] = avec; + tile[_kl*VW3 + 0][lid] = avec; #elif VW3 == 2 - tile[kl*VW3 + 0][lid] = avec.x; - tile[kl*VW3 + 1][lid] = avec.y; + tile[_kl*VW3 + 0][lid] = avec.x; + tile[_kl*VW3 + 1][lid] = avec.y; #elif VW3 == 4 - tile[kl*VW3 + 0][lid] = avec.x; - tile[kl*VW3 + 1][lid] = avec.y; - tile[kl*VW3 + 2][lid] = avec.z; - tile[kl*VW3 + 3][lid] = avec.w; + tile[_kl*VW3 + 0][lid] = avec.x; + tile[_kl*VW3 + 1][lid] = avec.y; + tile[_kl*VW3 + 2][lid] = avec.z; + tile[_kl*VW3 + 3][lid] = avec.w; #elif VW3 == 8 - tile[kl*VW3 + 0][lid] = avec.s0; - tile[kl*VW3 + 1][lid] = avec.s1; - tile[kl*VW3 + 2][lid] = avec.s2; - tile[kl*VW3 + 3][lid] = avec.s3; - tile[kl*VW3 + 4][lid] = avec.s4; - tile[kl*VW3 + 5][lid] = avec.s5; - tile[kl*VW3 + 6][lid] = avec.s6; - tile[kl*VW3 + 7][lid] = avec.s7; + tile[_kl*VW3 + 0][lid] = avec.s0; + tile[_kl*VW3 + 1][lid] = avec.s1; + tile[_kl*VW3 + 2][lid] = avec.s2; + tile[_kl*VW3 + 3][lid] = avec.s3; + tile[_kl*VW3 + 4][lid] = avec.s4; + tile[_kl*VW3 + 5][lid] = avec.s5; + tile[_kl*VW3 + 6][lid] = avec.s6; + tile[_kl*VW3 + 7][lid] = avec.s7; #elif VW3 == 16 - tile[kl*VW3 + 0][lid] = avec.s0; - tile[kl*VW3 + 1][lid] = avec.s1; - tile[kl*VW3 + 2][lid] = avec.s2; - tile[kl*VW3 + 3][lid] = avec.s3; - tile[kl*VW3 + 4][lid] = avec.s4; - tile[kl*VW3 + 5][lid] = avec.s5; - tile[kl*VW3 + 6][lid] = avec.s6; - tile[kl*VW3 + 7][lid] = avec.s7; - tile[kl*VW3 + 8][lid] = avec.s8; - tile[kl*VW3 + 9][lid] = avec.s9; - tile[kl*VW3 + 10][lid] = avec.sA; - tile[kl*VW3 + 11][lid] = avec.sB; - tile[kl*VW3 + 12][lid] = avec.sC; - tile[kl*VW3 + 13][lid] = avec.sD; - tile[kl*VW3 + 14][lid] = avec.sE; - tile[kl*VW3 + 15][lid] = avec.sF; + tile[_kl*VW3 + 0][lid] = avec.s0; + tile[_kl*VW3 + 1][lid] = avec.s1; + tile[_kl*VW3 + 2][lid] = avec.s2; + tile[_kl*VW3 + 3][lid] = avec.s3; + tile[_kl*VW3 + 4][lid] = avec.s4; + tile[_kl*VW3 + 5][lid] = avec.s5; + tile[_kl*VW3 + 6][lid] = avec.s6; + tile[_kl*VW3 + 7][lid] = avec.s7; + tile[_kl*VW3 + 8][lid] = avec.s8; + tile[_kl*VW3 + 9][lid] = avec.s9; + tile[_kl*VW3 + 10][lid] = avec.sA; + tile[_kl*VW3 + 11][lid] = avec.sB; + tile[_kl*VW3 + 12][lid] = avec.sC; + tile[_kl*VW3 + 13][lid] = avec.sD; + tile[_kl*VW3 + 14][lid] = avec.sE; + tile[_kl*VW3 + 15][lid] = avec.sF; #endif } @@ -274,11 +275,11 @@ void XgemvFastRot(const int m, const int n, // The multiply-add function (rotated) #pragma unroll - for (int kl=0; kl<WPT3/VW3; ++kl) { + for (int _kl = 0; _kl < WPT3/VW3; _kl += 1) { #pragma unroll - for (int v=0; v<VW3; ++v) { - real aval = tile[lid_mod*VW3 + v][lid_div * (WPT3/VW3) + kl]; - real xval = xlm[kl*VW3 + v]; + 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); } } diff --git a/src/kernels/level2/xger.opencl b/src/kernels/level2/xger.opencl index 1b9ded12..ca6071cd 100644 --- a/src/kernels/level2/xger.opencl +++ b/src/kernels/level2/xger.opencl @@ -28,7 +28,9 @@ void Xger(const int max1, const int max2, const real alpha = GetRealArg(arg_alpha); // Register storage for X and Y + #pragma promote_to_registers real xvalues[WPT]; + #pragma promote_to_registers real yvalues[WPT]; // Row-major version @@ -36,31 +38,31 @@ void Xger(const int max1, const int max2, // Loads the X-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id2 = w*get_global_size(1) + get_global_id(1); - xvalues[w] = LoadVector(id2, max2, xgm, x_offset, x_inc, false); + for (int _w = 0; _w < WPT; _w += 1) { + const int id2 = _w*get_global_size(1) + get_global_id(1); + xvalues[_w] = LoadVector(id2, max2, xgm, x_offset, x_inc, false); } // Loads the Y-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id1 = w*get_global_size(0) + get_global_id(0); - yvalues[w] = LoadVector(id1, max1, ygm, y_offset, y_inc, true); + for (int _w = 0; _w < WPT; _w += 1) { + const int id1 = _w*get_global_size(0) + get_global_id(0); + yvalues[_w] = LoadVector(id1, max1, ygm, y_offset, y_inc, true); } // Loops over the work per thread twice #pragma unroll - for (int w1=0; w1<WPT; ++w1) { + for (int _w1 = 0; _w1 < WPT; _w1 += 1) { #pragma unroll - for (int w2=0; w2<WPT; ++w2) { + for (int _w2 = 0; _w2 < WPT; _w2 += 1) { // Global thread IDs - const int id1 = w1*get_global_size(0) + get_global_id(0); - const int id2 = w2*get_global_size(1) + get_global_id(1); + const int id1 = _w1*get_global_size(0) + get_global_id(0); + const int id2 = _w2*get_global_size(1) + get_global_id(1); // Loads A, performs the operation, and stores the result into A MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld, - alpha, xvalues[w2], yvalues[w1], false); + alpha, xvalues[_w2], yvalues[_w1], false); } } } @@ -70,31 +72,31 @@ void Xger(const int max1, const int max2, // Loads the X-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id1 = w*get_global_size(0) + get_global_id(0); - xvalues[w] = LoadVector(id1, max1, xgm, x_offset, x_inc, false); + for (int _w = 0; _w < WPT; _w += 1) { + const int id1 = _w*get_global_size(0) + get_global_id(0); + xvalues[_w] = LoadVector(id1, max1, xgm, x_offset, x_inc, false); } // Loads the Y-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id2 = w*get_global_size(1) + get_global_id(1); - yvalues[w] = LoadVector(id2, max2, ygm, y_offset, y_inc, true); + for (int _w = 0; _w < WPT; _w += 1) { + const int id2 = _w*get_global_size(1) + get_global_id(1); + yvalues[_w] = LoadVector(id2, max2, ygm, y_offset, y_inc, true); } // Loops over the work per thread twice #pragma unroll - for (int w1=0; w1<WPT; ++w1) { + for (int _w1 = 0; _w1 < WPT; _w1 += 1) { #pragma unroll - for (int w2=0; w2<WPT; ++w2) { + for (int _w2 = 0; _w2 < WPT; _w2 += 1) { // Global thread IDs - const int id1 = w1*get_global_size(0) + get_global_id(0); - const int id2 = w2*get_global_size(1) + get_global_id(1); + const int id1 = _w1*get_global_size(0) + get_global_id(0); + const int id2 = _w2*get_global_size(1) + get_global_id(1); // Loads A, performs the operation, and stores the result into A MatrixUpdate(id1, id2, max1, max2, agm, a_offset, a_ld, - alpha, xvalues[w1], yvalues[w2], false); + alpha, xvalues[_w1], yvalues[_w2], false); } } } diff --git a/src/kernels/level3/copy_fast.opencl b/src/kernels/level3/copy_fast.opencl index 695b9003..ef8a9017 100644 --- a/src/kernels/level3/copy_fast.opencl +++ b/src/kernels/level3/copy_fast.opencl @@ -42,9 +42,9 @@ void CopyMatrixFast(const int ld, const real_arg arg_alpha) { const real alpha = GetRealArg(arg_alpha); #pragma unroll - for (int w_one=0; w_one<COPY_WPT; ++w_one) { + for (int _w_one = 0; _w_one < COPY_WPT; _w_one += 1) { const int id_one = get_global_id(0); - const int id_two = (get_group_id(1)*COPY_WPT + w_one) * COPY_DIMY + get_local_id(1); + const int id_two = (get_group_id(1)*COPY_WPT + _w_one) * COPY_DIMY + get_local_id(1); const int id = id_two*(ld/COPY_VW) + id_one; realC result; #if COPY_VW == 1 diff --git a/src/kernels/level3/copy_pad.opencl b/src/kernels/level3/copy_pad.opencl index 6eeadbd1..2e191514 100644 --- a/src/kernels/level3/copy_pad.opencl +++ b/src/kernels/level3/copy_pad.opencl @@ -35,11 +35,11 @@ INLINE_FUNC void _CopyPadMatrix(const int src_one, const int src_two, // Loops over the work per thread in both dimensions #pragma unroll - for (int w_one=0; w_one<PAD_WPTX; ++w_one) { - const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0); + for (int _w_one = 0; _w_one < PAD_WPTX; _w_one += 1) { + const int id_one = (get_group_id(0)*PAD_WPTX + _w_one) * PAD_DIMX + get_local_id(0); #pragma unroll - for (int w_two=0; w_two<PAD_WPTY; ++w_two) { - const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1); + for (int _w_two = 0; _w_two < PAD_WPTY; _w_two += 1) { + const int id_two = (get_group_id(1)*PAD_WPTY + _w_two) * PAD_DIMY + get_local_id(1); if (id_two < dest_two && id_one < dest_one) { // Loads data if the thread IDs are within bounds of the source matrix. Otherwise, set the @@ -91,11 +91,11 @@ INLINE_FUNC void _CopyMatrix(const int src_one, const int src_two, // Loops over the work per thread in both dimensions #pragma unroll - for (int w_one=0; w_one<PAD_WPTX; ++w_one) { - const int id_one = (get_group_id(0)*PAD_WPTX + w_one) * PAD_DIMX + get_local_id(0); + for (int _w_one = 0; _w_one < PAD_WPTX; _w_one += 1) { + const int id_one = (get_group_id(0)*PAD_WPTX + _w_one) * PAD_DIMX + get_local_id(0); #pragma unroll - for (int w_two=0; w_two<PAD_WPTY; ++w_two) { - const int id_two = (get_group_id(1)*PAD_WPTY + w_two) * PAD_DIMY + get_local_id(1); + for (int _w_two = 0; _w_two < PAD_WPTY; _w_two += 1) { + const int id_two = (get_group_id(1)*PAD_WPTY + _w_two) * PAD_DIMY + get_local_id(1); // Masking in case of triangular matrices: updates only the upper or lower part bool condition = true; |