summaryrefslogtreecommitdiff
path: root/src/kernels/level2/xher.opencl
diff options
context:
space:
mode:
authorCedric Nugteren <web@cedricnugteren.nl>2017-12-09 20:44:21 +0100
committerCedric Nugteren <web@cedricnugteren.nl>2017-12-09 20:44:21 +0100
commit9f02fb542ca659bf58d1efefdb334ea386ef10e8 (patch)
tree8669e215f38e8b00a4ee73a436479207f2afdbf5 /src/kernels/level2/xher.opencl
parentca5dbcd2bd31fb0a0e3f6c2f81b3c0fff6250738 (diff)
Completed kernel modifications for pre-processor of all other kernels
Diffstat (limited to 'src/kernels/level2/xher.opencl')
-rw-r--r--src/kernels/level2/xher.opencl24
1 files changed, 13 insertions, 11 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);
}
}
}