summaryrefslogtreecommitdiff
path: root/src/kernels/level3
diff options
context:
space:
mode:
Diffstat (limited to 'src/kernels/level3')
-rw-r--r--src/kernels/level3/copy_fast.opencl44
-rw-r--r--src/kernels/level3/copy_pad.opencl8
-rw-r--r--src/kernels/level3/transpose_fast.opencl46
-rw-r--r--src/kernels/level3/transpose_pad.opencl8
4 files changed, 97 insertions, 9 deletions
diff --git a/src/kernels/level3/copy_fast.opencl b/src/kernels/level3/copy_fast.opencl
index bfbfacd4..09e54e6d 100644
--- a/src/kernels/level3/copy_fast.opencl
+++ b/src/kernels/level3/copy_fast.opencl
@@ -38,13 +38,53 @@ R"(
__attribute__((reqd_work_group_size(COPY_DIMX, COPY_DIMY, 1)))
__kernel void CopyMatrixFast(const int ld,
__global const realC* restrict src,
- __global realC* dest) {
+ __global realC* dest,
+ const __constant real* restrict arg_alpha) {
+ const real alpha = arg_alpha[0];
#pragma unroll
for (int w_one=0; w_one<COPY_WPT; ++w_one) {
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 = id_two*(ld/COPY_VW) + id_one;
- dest[id] = src[id];
+ realC result;
+ #if COPY_VW == 1
+ Multiply(result, alpha, src[id]);
+ #elif COPY_VW == 2
+ Multiply(result.x, alpha, src[id].x);
+ Multiply(result.y, alpha, src[id].y);
+ #elif COPY_VW == 4
+ Multiply(result.x, alpha, src[id].x);
+ Multiply(result.y, alpha, src[id].y);
+ Multiply(result.z, alpha, src[id].z);
+ Multiply(result.w, alpha, src[id].w);
+ #elif COPY_VW == 8
+ Multiply(result.s0, alpha, src[id].s0);
+ Multiply(result.s1, alpha, src[id].s1);
+ Multiply(result.s2, alpha, src[id].s2);
+ Multiply(result.s3, alpha, src[id].s3);
+ Multiply(result.s4, alpha, src[id].s4);
+ Multiply(result.s5, alpha, src[id].s5);
+ Multiply(result.s6, alpha, src[id].s6);
+ Multiply(result.s7, alpha, src[id].s7);
+ #elif COPY_VW == 16
+ Multiply(result.s0, alpha, src[id].s0);
+ Multiply(result.s1, alpha, src[id].s1);
+ Multiply(result.s2, alpha, src[id].s2);
+ Multiply(result.s3, alpha, src[id].s3);
+ Multiply(result.s4, alpha, src[id].s4);
+ Multiply(result.s5, alpha, src[id].s5);
+ Multiply(result.s6, alpha, src[id].s6);
+ Multiply(result.s7, alpha, src[id].s7);
+ Multiply(result.s8, alpha, src[id].s8);
+ Multiply(result.s9, alpha, src[id].s9);
+ Multiply(result.sA, alpha, src[id].sA);
+ Multiply(result.sB, alpha, src[id].sB);
+ Multiply(result.sC, alpha, src[id].sC);
+ Multiply(result.sD, alpha, src[id].sD);
+ Multiply(result.sE, alpha, src[id].sE);
+ Multiply(result.sF, alpha, src[id].sF);
+ #endif
+ dest[id] = result;;
}
}
diff --git a/src/kernels/level3/copy_pad.opencl b/src/kernels/level3/copy_pad.opencl
index f211af0f..d276cc60 100644
--- a/src/kernels/level3/copy_pad.opencl
+++ b/src/kernels/level3/copy_pad.opencl
@@ -31,7 +31,9 @@ __kernel void CopyPadMatrix(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];
// Loops over the work per thread in both dimensions
#pragma unroll
@@ -52,7 +54,7 @@ __kernel void CopyPadMatrix(const int src_one, const int src_two,
// Stores the value in the destination matrix
if (do_conjugate == 1) { COMPLEX_CONJUGATE(value); }
- dest[id_two*dest_ld + id_one + dest_offset] = value;
+ Multiply(dest[id_two*dest_ld + id_one + dest_offset], alpha, value);
}
}
}
@@ -70,8 +72,10 @@ __kernel void CopyMatrix(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];
// Loops over the work per thread in both dimensions
#pragma unroll
@@ -94,7 +98,7 @@ __kernel void CopyMatrix(const int src_one, const int src_two,
if (id_two < dest_two && id_one < dest_one) {
real value = src[id_two*src_ld + id_one + src_offset];
if (diagonal_imag_zero == 1 && id_one == id_two) { ImagToZero(value); }
- dest[id_two*dest_ld + id_one + dest_offset] = value;
+ Multiply(dest[id_two*dest_ld + id_one + dest_offset], alpha, value);
}
}
}
diff --git a/src/kernels/level3/transpose_fast.opencl b/src/kernels/level3/transpose_fast.opencl
index 08266461..d5c46a30 100644
--- a/src/kernels/level3/transpose_fast.opencl
+++ b/src/kernels/level3/transpose_fast.opencl
@@ -39,7 +39,9 @@ R"(
__attribute__((reqd_work_group_size(TRA_DIM, TRA_DIM, 1)))
__kernel void TransposeMatrixFast(const int ld,
__global const realT* restrict src,
- __global realT* dest) {
+ __global realT* dest,
+ const __constant real* restrict arg_alpha) {
+ const real alpha = arg_alpha[0];
// Sets the group identifiers. They might be 'shuffled' around to distribute work in a different
// way over workgroups, breaking memory-bank dependencies.
@@ -117,12 +119,50 @@ __kernel void TransposeMatrixFast(const int ld,
results[15] = (realT) {v[0].sF, v[1].sF, v[2].sF, v[3].sF, v[4].sF, v[5].sF, v[6].sF, v[7].sF, v[8].sF, v[9].sF, v[10].sF, v[11].sF, v[12].sF, v[13].sF, v[14].sF, v[15].sF};
#endif
- // Stores the results into the destination matrix
+ // Multiplies by alpha and then stores the results into the destination matrix
#pragma unroll
for (int w_two=0; w_two<TRA_WPT; ++w_two) {
+ realT result;
+ #if TRA_WPT == 1
+ Multiply(result, alpha, results[w_two]);
+ #elif TRA_WPT == 2
+ Multiply(result.x, alpha, results[w_two].x);
+ Multiply(result.y, alpha, results[w_two].y);
+ #elif TRA_WPT == 4
+ Multiply(result.x, alpha, results[w_two].x);
+ Multiply(result.y, alpha, results[w_two].y);
+ Multiply(result.z, alpha, results[w_two].z);
+ Multiply(result.w, alpha, results[w_two].w);
+ #elif TRA_WPT == 8
+ Multiply(result.s0, alpha, results[w_two].s0);
+ Multiply(result.s1, alpha, results[w_two].s1);
+ Multiply(result.s2, alpha, results[w_two].s2);
+ Multiply(result.s3, alpha, results[w_two].s3);
+ Multiply(result.s4, alpha, results[w_two].s4);
+ Multiply(result.s5, alpha, results[w_two].s5);
+ Multiply(result.s6, alpha, results[w_two].s6);
+ Multiply(result.s7, alpha, results[w_two].s7);
+ #elif TRA_WPT == 16
+ Multiply(result.s0, alpha, results[w_two].s0);
+ Multiply(result.s1, alpha, results[w_two].s1);
+ Multiply(result.s2, alpha, results[w_two].s2);
+ Multiply(result.s3, alpha, results[w_two].s3);
+ Multiply(result.s4, alpha, results[w_two].s4);
+ Multiply(result.s5, alpha, results[w_two].s5);
+ Multiply(result.s6, alpha, results[w_two].s6);
+ Multiply(result.s7, alpha, results[w_two].s7);
+ Multiply(result.s8, alpha, results[w_two].s8);
+ Multiply(result.s9, alpha, results[w_two].s9);
+ Multiply(result.sA, alpha, results[w_two].sA);
+ Multiply(result.sB, alpha, results[w_two].sB);
+ Multiply(result.sC, alpha, results[w_two].sC);
+ Multiply(result.sD, alpha, results[w_two].sD);
+ Multiply(result.sE, alpha, results[w_two].sE);
+ Multiply(result.sF, alpha, results[w_two].sF);
+ #endif
const int id_one = gid0*TRA_DIM + get_local_id(0);
const int id_two = (gid1*TRA_DIM + get_local_id(1))*TRA_WPT + w_two;
- dest[id_two*(ld/TRA_WPT) + id_one] = results[w_two];
+ dest[id_two*(ld/TRA_WPT) + id_one] = result;
}
}
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);
}
}
}
generated by cgit v1.2.3 (git 2.39.1) at 2024-07-07 06:45:29 +0000