summaryrefslogtreecommitdiff
path: root/src/kernels/level3/transpose_pad.opencl
diff options
context:
space:
mode:
Diffstat (limited to 'src/kernels/level3/transpose_pad.opencl')
-rw-r--r--src/kernels/level3/transpose_pad.opencl8
1 files changed, 6 insertions, 2 deletions
diff --git a/src/kernels/level3/transpose_pad.opencl b/src/kernels/level3/transpose_pad.opencl
index 38c23346..2de0c7bd 100644
--- a/src/kernels/level3/transpose_pad.opencl
+++ b/src/kernels/level3/transpose_pad.opencl
@@ -31,7 +31,9 @@ __kernel void TransposePadMatrix(const int src_one, const int src_two,
const int dest_one, const int dest_two,
const int dest_ld, const int dest_offset,
__global real* dest,
+ const __constant real* restrict arg_alpha,
const int do_conjugate) {
+ const real alpha = arg_alpha[0];
// Local memory to store a tile of the matrix (for coalescing)
__local real tile[PADTRA_WPT*PADTRA_TILE][PADTRA_WPT*PADTRA_TILE + PADTRA_PAD];
@@ -75,7 +77,7 @@ __kernel void TransposePadMatrix(const int src_one, const int src_two,
if ((id_dest_one < dest_one) && (id_dest_two < dest_two)) {
real value = tile[get_local_id(0)*PADTRA_WPT + w_two][get_local_id(1)*PADTRA_WPT + w_one];
if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
- dest[id_dest_two*dest_ld + id_dest_one + dest_offset] = value;
+ Multiply(dest[id_dest_two*dest_ld + id_dest_one + dest_offset], alpha, value);
}
}
}
@@ -93,8 +95,10 @@ __kernel void TransposeMatrix(const int src_one, const int src_two,
const int dest_one, const int dest_two,
const int dest_ld, const int dest_offset,
__global real* dest,
+ const __constant real* restrict arg_alpha,
const int upper, const int lower,
const int diagonal_imag_zero) {
+ const real alpha = arg_alpha[0];
// Local memory to store a tile of the matrix (for coalescing)
__local real tile[PADTRA_WPT*PADTRA_TILE][PADTRA_WPT*PADTRA_TILE + PADTRA_PAD];
@@ -143,7 +147,7 @@ __kernel void TransposeMatrix(const int src_one, const int src_two,
if ((id_dest_one < dest_one) && (id_dest_two < dest_two)) {
real value = tile[get_local_id(0)*PADTRA_WPT + w_two][get_local_id(1)*PADTRA_WPT + w_one];
if (diagonal_imag_zero == 1 && id_dest_one == id_dest_two) { ImagToZero(value); }
- dest[id_dest_two*dest_ld + id_dest_one + dest_offset] = value;
+ Multiply(dest[id_dest_two*dest_ld + id_dest_one + dest_offset], alpha, value);
}
}
}