diff options
| author | Cedric Nugteren <web@cedricnugteren.nl> | 2016-07-10 20:32:01 +0200 |
|---|---|---|
| committer | Cedric Nugteren <web@cedricnugteren.nl> | 2016-07-10 20:32:01 +0200 |
| commit | c87e877bf23d2fe38a7da2898e1734a3cdeaf48c (patch) | |
| tree | d091dfdd826dd11e5c9e533eb46b22aeb7f6f823 /src/kernels/level3/xgemm_part2.opencl | |
| parent | 57f09178d89a1cf4f38a0bb338c864ed850d5470 (diff) | |
Now passing alpha/beta to the kernel as arguments as before fp16 support; in case of fp16 arguments are cast on host and in kernel
Diffstat (limited to 'src/kernels/level3/xgemm_part2.opencl')
| -rw-r--r-- | src/kernels/level3/xgemm_part2.opencl | 24 |
1 files changed, 12 insertions, 12 deletions
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 |