cl_khr_expect_assume Extension Support (CNugteren#610)

* Added support for cl_khr_expect_assume extension to enable some optimizations by the compiler.

* Undid 'use' of parameters with value 0 in gemv.

---------

Co-authored-by: Cedric Nugteren <web@cedricnugteren.nl>

Author: EJainDev

src/kernels/common.opencl

@@ -12,6 +12,14 @@
 R"(
 // =================================================================================================
 
+#if defined(cl_khr_expect_assume)
+#pragma OPENCL EXTENSION cl_khr_expect_assume : enable
+#endif
+
+#if !defined(__has_builtin)
+#define __has_builtin(x) 0
+#endif
+
 // Parameters set by the tuner or by the database. Here they are given a basic default value in case
 // this file is used outside of the CLBlast library.
 #ifndef PRECISION

src/kernels/level1/xaxpy.opencl

@@ -45,6 +45,11 @@ void Xaxpy(const int n, const real_arg arg_alpha,
 void XaxpyFaster(const int n, const real_arg arg_alpha,
                  const __global realV* restrict xgm,
                  __global realV* ygm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+#endif
+
   const real alpha = GetRealArg(arg_alpha);
 
   const int num_usefull_threads = n / (VW * WPT);
@@ -69,6 +74,12 @@ void XaxpyFaster(const int n, const real_arg arg_alpha,
 void XaxpyFastest(const int n, const real_arg arg_alpha,
                   const __global realV* restrict xgm,
                   __global realV* ygm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+  __builtin_assume(n % WGS == 0);
+#endif
+
   const real alpha = GetRealArg(arg_alpha);
 
   #pragma unroll

src/kernels/level1/xcopy.opencl

@@ -45,6 +45,11 @@ void Xcopy(const int n,
 void XcopyFast(const int n,
                const __global realV* restrict xgm,
                __global realV* ygm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+  __builtin_assume(n % WGS == 0);
+#endif
   #pragma unroll
   for (int _w = 0; _w < WPT; _w += 1) {
     const int id = _w*get_global_size(0) + get_global_id(0);

src/kernels/level1/xhad.opencl

@@ -98,6 +98,10 @@ void Xhad(const int n, const real_arg arg_alpha, const real_arg arg_beta,
 void XhadFaster(const int n, const real_arg arg_alpha, const real_arg arg_beta,
                 const __global realV* restrict xgm, const __global realV* restrict ygm,
                 __global realV* zgm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+#endif
   const real alpha = GetRealArg(arg_alpha);
   const real beta = GetRealArg(arg_beta);
 
@@ -129,6 +133,11 @@ void XhadFaster(const int n, const real_arg arg_alpha, const real_arg arg_beta,
 void XhadFastest(const int n, const real_arg arg_alpha, const real_arg arg_beta,
                  const __global realV* restrict xgm, const __global realV* restrict ygm,
                  __global realV* zgm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+  __builtin_assume(n % WGS == 0);
+#endif
   const real alpha = GetRealArg(arg_alpha);
   const real beta = GetRealArg(arg_beta);
 

src/kernels/level1/xscal.opencl

@@ -47,6 +47,11 @@ void Xscal(const int n, const real_arg arg_alpha,
 #endif
 void XscalFast(const int n, const real_arg arg_alpha,
                __global realV* xgm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+  __builtin_assume(n % WGS == 0);
+#endif
   const real alpha = GetRealArg(arg_alpha);
 
   #pragma unroll

src/kernels/level1/xswap.opencl

@@ -47,6 +47,11 @@ void Xswap(const int n,
 void XswapFast(const int n,
                __global realV* xgm,
                __global realV* ygm) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(n % VW == 0);
+  __builtin_assume(n % WPT == 0);
+  __builtin_assume(n % WGS == 0);
+#endif
   #pragma unroll
   for (int _w = 0; _w < WPT; _w += 1) {
     const int id = _w*get_global_size(0) + get_global_id(0);

src/kernels/level2/xgemv_fast.opencl

@@ -98,6 +98,11 @@ void XgemvFast(const int m, const int n,
                __global real* ygm, const int y_offset, const int y_inc,
                const int do_conjugate, const int parameter,
                const int kl_unused, const int ku_unused) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(m % WGS2 == 0);
+  __builtin_assume(n % WGS2 == 0);
+  __builtin_assume(a_ld % VW2 == 0);
+#endif
   const real alpha = GetRealArg(arg_alpha);
   const real beta = GetRealArg(arg_beta);
 
@@ -205,6 +210,12 @@ void XgemvFastRot(const int m, const int n,
                   __global real* ygm, const int y_offset, const int y_inc,
                   const int do_conjugate, const int parameter,
                   const int kl_unused, const int ku_unused) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(m % WGS3 == 0);
+  __builtin_assume(n % WGS3 == 0);
+  __builtin_assume(a_ld % VW3 == 0);
+#endif
+
   const real alpha = GetRealArg(arg_alpha);
   const real beta = GetRealArg(arg_beta);
 

src/kernels/level3/copy_fast.opencl

@@ -40,6 +40,10 @@ void CopyMatrixFast(const int ld,
                     __global const realC* restrict src,
                     __global realC* dest,
                     const real_arg arg_alpha) {
+#if __has_builtin(__builtin_assume)
+  __builtin_assume(ld % COPY_VW == 0);
+#endif
+
   const real alpha = GetRealArg(arg_alpha);
   #pragma unroll
   for (int _w_one = 0; _w_one < COPY_WPT; _w_one += 1) {