diff options
Diffstat (limited to 'src/kernels/level2')
-rw-r--r-- | src/kernels/level2/xher.opencl | 24 | ||||
-rw-r--r-- | src/kernels/level2/xher2.opencl | 40 |
2 files changed, 35 insertions, 29 deletions
diff --git a/src/kernels/level2/xher.opencl b/src/kernels/level2/xher.opencl index b0772218..8a57bdfc 100644 --- a/src/kernels/level2/xher.opencl +++ b/src/kernels/level2/xher.opencl @@ -27,32 +27,34 @@ void Xher(const int n, const real alpha = GetRealArg(arg_alpha); // Register storage for X and XT + #pragma promote_to_registers real xvalues[WPT]; + #pragma promote_to_registers real xtvalues[WPT]; // 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, n, xgm, x_offset, x_inc, !is_rowmajor); + for (int _w = 0; _w < WPT; _w += 1) { + const int id2 = _w*get_global_size(1) + get_global_id(1); + xvalues[_w] = LoadVector(id2, n, xgm, x_offset, x_inc, !is_rowmajor); } // Loads the X-transposed-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id1 = w*get_global_size(0) + get_global_id(0); - xtvalues[w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor); + for (int _w = 0; _w < WPT; _w += 1) { + const int id1 = _w*get_global_size(0) + get_global_id(0); + xtvalues[_w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor); } // 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); // Skip these threads if they do not contain threads contributing to the matrix-triangle if ((is_upper && (id1 > id2)) || (!is_upper && (id2 > id1))) { @@ -61,7 +63,7 @@ void Xher(const int n, // Loads A, performs the operation, and stores the result into A else { - MatrixUpdate(id1, id2, n, n, agm, a_offset, a_ld, alpha, xvalues[w2], xtvalues[w1], is_upper); + MatrixUpdate(id1, id2, n, n, agm, a_offset, a_ld, alpha, xvalues[_w2], xtvalues[_w1], is_upper); } } } diff --git a/src/kernels/level2/xher2.opencl b/src/kernels/level2/xher2.opencl index 00a756c9..73305149 100644 --- a/src/kernels/level2/xher2.opencl +++ b/src/kernels/level2/xher2.opencl @@ -28,37 +28,41 @@ void Xher2(const int n, 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]; + #pragma promote_to_registers real xtvalues[WPT]; + #pragma promote_to_registers real ytvalues[WPT]; // 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, n, xgm, x_offset, x_inc, !is_rowmajor); + for (int _w = 0; _w < WPT; _w += 1) { + const int id2 = _w*get_global_size(1) + get_global_id(1); + xvalues[_w] = LoadVector(id2, n, xgm, x_offset, x_inc, !is_rowmajor); } // Loads the X-transposed-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id1 = w*get_global_size(0) + get_global_id(0); - xtvalues[w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor); + for (int _w = 0; _w < WPT; _w += 1) { + const int id1 = _w*get_global_size(0) + get_global_id(0); + xtvalues[_w] = LoadVector(id1, n, xgm, x_offset, x_inc, is_rowmajor); } // 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, n, ygm, y_offset, y_inc, is_rowmajor); + for (int _w = 0; _w < WPT; _w += 1) { + const int id1 = _w*get_global_size(0) + get_global_id(0); + yvalues[_w] = LoadVector(id1, n, ygm, y_offset, y_inc, is_rowmajor); } // Loads the Y-transposed-vector #pragma unroll - for (int w=0; w<WPT; ++w) { - const int id2 = w*get_global_size(1) + get_global_id(1); - ytvalues[w] = LoadVector(id2, n, ygm, y_offset, y_inc, !is_rowmajor); + for (int _w = 0; _w < WPT; _w += 1) { + const int id2 = _w*get_global_size(1) + get_global_id(1); + ytvalues[_w] = LoadVector(id2, n, ygm, y_offset, y_inc, !is_rowmajor); } // Sets the proper value of alpha in case conjugation is needed @@ -75,13 +79,13 @@ void Xher2(const int n, // 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); // Skip these threads if they do not contain threads contributing to the matrix-triangle if ((is_upper && (id1 > id2)) || (!is_upper && (id2 > id1))) { @@ -91,8 +95,8 @@ void Xher2(const int n, // Loads A, performs the operation, and stores the result into A else { MatrixUpdate2(id1, id2, n, n, agm, a_offset, a_ld, - alpha1, xvalues[w2], yvalues[w1], - alpha2, xtvalues[w1], ytvalues[w2], is_upper); + alpha1, xvalues[_w2], yvalues[_w1], + alpha2, xtvalues[_w1], ytvalues[_w2], is_upper); } } } |