diff options
author | Cedric Nugteren <web@cedricnugteren.nl> | 2016-05-16 12:37:24 +0200 |
---|---|---|
committer | Cedric Nugteren <web@cedricnugteren.nl> | 2016-05-16 12:37:24 +0200 |
commit | af2ac6221288ec101a69018e960bb004ad698efe (patch) | |
tree | 9173d204820075598dd2fe13a0cb1771fd732945 /src/kernels/level3/xgemm_part2.opencl | |
parent | 591e343ec94077f873b1aa12052a4ce55ae80200 (diff) |
Prepared GEMM and supporting kernels and tuners for half-precision support
Diffstat (limited to 'src/kernels/level3/xgemm_part2.opencl')
-rw-r--r-- | src/kernels/level3/xgemm_part2.opencl | 15 |
1 files changed, 12 insertions, 3 deletions
diff --git a/src/kernels/level3/xgemm_part2.opencl b/src/kernels/level3/xgemm_part2.opencl index c0760db6..a8c8ebf5 100644 --- a/src/kernels/level3/xgemm_part2.opencl +++ b/src/kernels/level3/xgemm_part2.opencl @@ -263,10 +263,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 real alpha, const real beta, + const __constant real* restrict arg_alpha, + const __constant real* restrict 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]; // Skip these threads if they do not contain threads contributing to the upper-triangle if (get_group_id(1)*NWG < get_group_id(0)*MWG) { @@ -300,10 +303,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 real alpha, const real beta, + const __constant real* restrict arg_alpha, + const __constant real* restrict 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]; // Skip these threads if they do not contain threads contributing to the lower-triangle if (get_group_id(1)*NWG > get_group_id(0)*MWG) { @@ -341,10 +347,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 real alpha, const real beta, + const __constant real* restrict arg_alpha, + const __constant real* restrict 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]; // Allocates workgroup-private memory (local memory) #if SA == 1 |