Add cl_khr_fp64 when using double precision

Author: shehzan10

src/library/blas/AutoGemm/KernelOpenCL.py

@@ -25,6 +25,13 @@ def makeOpenCLKernelString(kernel):
   kStr += "/* %s */" % kernel.getName()
   kStr += endLine
 
+  ####################################
+  # Double precision pragma
+  prec = kernel.getName()[0].lower()
+  if prec == "d" or prec == "z":
+    kStr += endLine
+    kStr += "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" + endLine
+
   ####################################
   # kernel parameters
   kStr += endLine

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NN_B0_MX048_NX048_KX08_src.cpp

@@ -18,6 +18,7 @@ const unsigned int dgemm_Col_NN_B0_MX048_NX048_KX08_microTileNumCols = 6;
 const unsigned int dgemm_Col_NN_B0_MX048_NX048_KX08_unroll = 8;
 
 const char * const dgemm_Col_NN_B0_MX048_NX048_KX08_src = STRINGIFY(
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable \n
 
 #define  M6x6 \
             rA[0] = lA[offA + 0];\

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NN_B1_MX048_NX048_KX08_src.cpp

@@ -18,6 +18,7 @@ const unsigned int dgemm_Col_NN_B1_MX048_NX048_KX08_microTileNumCols = 6;
 const unsigned int dgemm_Col_NN_B1_MX048_NX048_KX08_unroll = 8;
 
 const char * const dgemm_Col_NN_B1_MX048_NX048_KX08_src = STRINGIFY(
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable   \n
 
 #define  M6x6 \
             rA[0] = lA[offA + 0];                       \

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NT_B0_MX048_NX048_KX08_src.cpp

@@ -18,6 +18,7 @@ const unsigned int dgemm_Col_NT_B0_MX048_NX048_KX08_microTileNumCols = 6;
 const unsigned int dgemm_Col_NT_B0_MX048_NX048_KX08_unroll = 8;
 
 const char * const dgemm_Col_NT_B0_MX048_NX048_KX08_src = STRINGIFY(
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable   \n
 \n
 \ntypedef union _GPtr {
 \n  __global float *f;

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_NT_B1_MX048_NX048_KX08_src.cpp

@@ -18,6 +18,7 @@ const unsigned int dgemm_Col_NT_B1_MX048_NX048_KX08_microTileNumCols = 6;
 const unsigned int dgemm_Col_NT_B1_MX048_NX048_KX08_unroll = 8;
 
 const char * const dgemm_Col_NT_B1_MX048_NX048_KX08_src = STRINGIFY(
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable   \n
 \n
 \ntypedef union _GPtr {
 \n  __global float *f;

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_TN_B0_MX048_NX048_KX08_src.cpp

@@ -18,6 +18,7 @@ const unsigned int dgemm_Col_TN_B0_MX048_NX048_KX08_microTileNumCols = 6;
 const unsigned int dgemm_Col_TN_B0_MX048_NX048_KX08_unroll = 8;
 
 const char * const dgemm_Col_TN_B0_MX048_NX048_KX08_src = STRINGIFY(
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable   \n
 
 __attribute__( (reqd_work_group_size(8, 8, 1)) )
 __kernel void dgemm_Col_TN_B0_MX048_NX048_KX08_src (

src/library/blas/AutoGemm/UserGemmKernelSources/dgemm_Col_TN_B1_MX048_NX048_KX08_src.cpp

@@ -18,6 +18,7 @@ const unsigned int dgemm_Col_TN_B1_MX048_NX048_KX08_microTileNumCols = 6;
 const unsigned int dgemm_Col_TN_B1_MX048_NX048_KX08_unroll = 8;
 
 const char * const dgemm_Col_TN_B1_MX048_NX048_KX08_src = STRINGIFY(
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable   \n
 
 __attribute__( (reqd_work_group_size(8, 8, 1)) )
 __kernel void dgemm_Col_TN_B1_MX048_NX048_KX08_src (