diff options
Diffstat (limited to 'src/kernels/level3')
-rw-r--r-- | src/kernels/level3/copy_fast.opencl | 4 | ||||
-rw-r--r-- | src/kernels/level3/copy_pad.opencl | 8 | ||||
-rw-r--r-- | src/kernels/level3/transpose_fast.opencl | 4 | ||||
-rw-r--r-- | src/kernels/level3/transpose_pad.opencl | 8 | ||||
-rw-r--r-- | src/kernels/level3/xgemm_part2.opencl | 24 |
5 files changed, 24 insertions, 24 deletions
diff --git a/src/kernels/level3/copy_fast.opencl b/src/kernels/level3/copy_fast.opencl index 09e54e6d..dd975bf1 100644 --- a/src/kernels/level3/copy_fast.opencl +++ b/src/kernels/level3/copy_fast.opencl @@ -39,8 +39,8 @@ __attribute__((reqd_work_group_size(COPY_DIMX, COPY_DIMY, 1))) __kernel void CopyMatrixFast(const int ld, __global const realC* restrict src, __global realC* dest, - const __constant real* restrict arg_alpha) { - const real alpha = arg_alpha[0]; + const real_arg arg_alpha) { + const real alpha = GetRealArg(arg_alpha); #pragma unroll for (int w_one=0; w_one<COPY_WPT; ++w_one) { const int id_one = get_global_id(0); diff --git a/src/kernels/level3/copy_pad.opencl b/src/kernels/level3/copy_pad.opencl index d276cc60..d0771c31 100644 --- a/src/kernels/level3/copy_pad.opencl +++ b/src/kernels/level3/copy_pad.opencl @@ -31,9 +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 real_arg arg_alpha, const int do_conjugate) { - const real alpha = arg_alpha[0]; + const real alpha = GetRealArg(arg_alpha); // Loops over the work per thread in both dimensions #pragma unroll @@ -72,10 +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 real_arg arg_alpha, const int upper, const int lower, const int diagonal_imag_zero) { - const real alpha = arg_alpha[0]; + const real alpha = GetRealArg(arg_alpha); // Loops over the work per thread in both dimensions #pragma unroll diff --git a/src/kernels/level3/transpose_fast.opencl b/src/kernels/level3/transpose_fast.opencl index d5c46a30..ea343533 100644 --- a/src/kernels/level3/transpose_fast.opencl +++ b/src/kernels/level3/transpose_fast.opencl @@ -40,8 +40,8 @@ __attribute__((reqd_work_group_size(TRA_DIM, TRA_DIM, 1))) __kernel void TransposeMatrixFast(const int ld, __global const realT* restrict src, __global realT* dest, - const __constant real* restrict arg_alpha) { - const real alpha = arg_alpha[0]; + const real_arg arg_alpha) { + const real alpha = GetRealArg(arg_alpha); // Sets the group identifiers. They might be 'shuffled' around to distribute work in a different // way over workgroups, breaking memory-bank dependencies. diff --git a/src/kernels/level3/transpose_pad.opencl b/src/kernels/level3/transpose_pad.opencl index 2de0c7bd..2e20d667 100644 --- a/src/kernels/level3/transpose_pad.opencl +++ b/src/kernels/level3/transpose_pad.opencl @@ -31,9 +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 real_arg arg_alpha, const int do_conjugate) { - const real alpha = arg_alpha[0]; + const real alpha = GetRealArg(arg_alpha); // Local memory to store a tile of the matrix (for coalescing) __local real tile[PADTRA_WPT*PADTRA_TILE][PADTRA_WPT*PADTRA_TILE + PADTRA_PAD]; @@ -95,10 +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 real_arg arg_alpha, const int upper, const int lower, const int diagonal_imag_zero) { - const real alpha = arg_alpha[0]; + const real alpha = GetRealArg(arg_alpha); // Local memory to store a tile of the matrix (for coalescing) __local real tile[PADTRA_WPT*PADTRA_TILE][PADTRA_WPT*PADTRA_TILE + PADTRA_PAD]; diff --git a/src/kernels/level3/xgemm_part2.opencl b/src/kernels/level3/xgemm_part2.opencl index 42c1127c..87e28cb5 100644 --- a/src/kernels/level3/xgemm_part2.opencl +++ b/src/kernels/level3/xgemm_part2.opencl @@ -270,13 +270,13 @@ inline void XgemmBody(const int kSizeM, const int kSizeN, const int kSizeK, // Main entry point of the kernel. This is the upper-triangular version. __attribute__((reqd_work_group_size(MDIMC, NDIMC, 1))) __kernel void XgemmUpper(const int kSizeN, const int kSizeK, - const __constant real* restrict arg_alpha, - const __constant real* restrict arg_beta, + const real_arg arg_alpha, + const real_arg arg_beta, const __global realM* restrict agm, const __global realN* restrict bgm, __global realM* cgm) { - const real alpha = arg_alpha[0]; - const real beta = arg_beta[0]; + const real alpha = GetRealArg(arg_alpha); + const real beta = GetRealArg(arg_beta); // Skip these threads if they do not contain threads contributing to the upper-triangle if (GetGroupID1()*NWG < GetGroupID0()*MWG) { @@ -310,13 +310,13 @@ __kernel void XgemmUpper(const int kSizeN, const int kSizeK, // Main entry point of the kernel. This is the lower-triangular version. __attribute__((reqd_work_group_size(MDIMC, NDIMC, 1))) __kernel void XgemmLower(const int kSizeN, const int kSizeK, - const __constant real* restrict arg_alpha, - const __constant real* restrict arg_beta, + const real_arg arg_alpha, + const real_arg arg_beta, const __global realM* restrict agm, const __global realN* restrict bgm, __global realM* cgm) { - const real alpha = arg_alpha[0]; - const real beta = arg_beta[0]; + const real alpha = GetRealArg(arg_alpha); + const real beta = GetRealArg(arg_beta); // Skip these threads if they do not contain threads contributing to the lower-triangle if (GetGroupID1()*NWG > GetGroupID0()*MWG) { @@ -354,13 +354,13 @@ __kernel void XgemmLower(const int kSizeN, const int kSizeK, // Main entry point of the kernel. This is the regular full version. __attribute__((reqd_work_group_size(MDIMC, NDIMC, 1))) __kernel void Xgemm(const int kSizeM, const int kSizeN, const int kSizeK, - const __constant real* restrict arg_alpha, - const __constant real* restrict arg_beta, + const real_arg arg_alpha, + const real_arg arg_beta, const __global realM* restrict agm, const __global realN* restrict bgm, __global realM* cgm) { - const real alpha = arg_alpha[0]; - const real beta = arg_beta[0]; + const real alpha = GetRealArg(arg_alpha); + const real beta = GetRealArg(arg_beta); // Allocates workgroup-private memory (local memory) #if SA == 1 |