Merge pull request #70 from CNugteren/development

Update to version 0.8.0

1 files changed, 15 insertions, 3 deletions

diff --git a/src/kernels/level3/xgemm_part2.opencl b/src/kernels/level3/xgemm_part2.opencl
index 599e01d5..42c1127c 100644
--- a/src/kernels/level3/xgemm_part2.opencl
+++ b/src/kernels/level3/xgemm_part2.opencl
@@ -258,6 +258,9 @@ inline void XgemmBody(const int kSizeM, const int kSizeN, const int kSizeK,
       barrier(CLK_LOCAL_MEM_FENCE);
     #endif
   }
+  #if GLOBAL_MEM_FENCE == 1
+    barrier(CLK_GLOBAL_MEM_FENCE);
+  #endif
 }
 
 // =================================================================================================
@@ -267,10 +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 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 (GetGroupID1()*NWG < GetGroupID0()*MWG) {
@@ -304,10 +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 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 (GetGroupID1()*NWG > GetGroupID0()*MWG) {
@@ -345,10 +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 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