diff options
author | Cedric Nugteren <web@cedricnugteren.nl> | 2017-11-29 20:21:08 +0100 |
---|---|---|
committer | Cedric Nugteren <web@cedricnugteren.nl> | 2017-11-29 20:21:08 +0100 |
commit | 93ffb876c60838bee75d3bb25ebbcbfce02e2cc7 (patch) | |
tree | 8578cb8ab62db14caa40d0f2647b8f06806b31fd /src/kernels/level2/xger.opencl | |
parent | 0dde6af703816adb0d53f00a88d007199c953042 (diff) |
Reformatted unrollable kernel loops and added the new promote_to_registers pragma for several kernels
Diffstat (limited to 'src/kernels/level2/xger.opencl')
-rw-r--r-- | src/kernels/level2/xger.opencl | 46 |
1 files changed, 24 insertions, 22 deletions
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); } } } |