Control: tags -1 patch
Running the clblas examples finds that at least cher, chpr and srotg are
also affected; the attached extends the patch to cover these.
It has *not* been tested beyond running the examples (which I suspect
don't check correctness): the test suite fails to build with what looks
like https://github.com/clMathLibraries/clBLAS/issues/338.
Description: Don't use double literals in single precision functions
Allows it to work on single-precision-only hardware/ICDs e.g. beignet
Author: Rebecca N. Palmer <[email protected]>
Bug-Debian: https://bugs.debian.org/877316
Forwarded: https://github.com/clMathLibraries/clBLAS/issues/321
--- a/src/library/blas/AutoGemm/KernelOpenCL.py
+++ b/src/library/blas/AutoGemm/KernelOpenCL.py
@@ -364,6 +364,8 @@ def makeOpenCLKernelString(kernel):
zeroString = "(float2)(0.f, 0.f)"
elif kernel.precision == "z":
zeroString = "(double2)(0.0, 0.0)"
+ elif kernel.precision == "s":
+ zeroString = "0.0f"
else:
zeroString = "0.0"
for a in range(0, int(numALoads)):
--- a/src/library/blas/AutoGemm/UserGemmKernelSources/sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_src.cpp
+++ b/src/library/blas/AutoGemm/UserGemmKernelSources/sgemm_Col_NN_B1_MX032_NX032_KX16_BRANCH_src.cpp
@@ -86,11 +86,11 @@ __kernel void sgemm_Col_NN_B1_MX032_NX03
__local float* plB = lB + idx*33+idy;
barrier(CLK_LOCAL_MEM_FENCE);
- plB[0] = CurrentOffSetB>=N?0.0:B[0];
- plB[16] = CurrentOffSetB+16>=N?0.0:B[16*ldb];
+ plB[0] = CurrentOffSetB>=N?0.0f:B[0];
+ plB[16] = CurrentOffSetB+16>=N?0.0f:B[16*ldb];
- plA[0] = CurrentOffSetA>=M?0.0:A[0];
- plA[16] = CurrentOffSetA+16>=M?0.0:A[16];
+ plA[0] = CurrentOffSetA>=M?0.0f:A[0];
+ plA[16] = CurrentOffSetA+16>=M?0.0f:A[16];
barrier(CLK_LOCAL_MEM_FENCE);
--- a/src/library/blas/AutoGemm/UserGemmKernelSources/sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_src.cpp
+++ b/src/library/blas/AutoGemm/UserGemmKernelSources/sgemm_Col_NT_B1_MX032_NX032_KX16_BRANCH_src.cpp
@@ -85,11 +85,11 @@ __kernel void sgemm_Col_NT_B1_MX032_NX03
__local float* plB = lB + idy*33+idx;
barrier(CLK_LOCAL_MEM_FENCE);
- plB[0] = CurrentOffSetB>=N?0.0:B[0];
- plB[16] = CurrentOffSetB+16>=N?0.0:B[16];
+ plB[0] = CurrentOffSetB>=N?0.0f:B[0];
+ plB[16] = CurrentOffSetB+16>=N?0.0f:B[16];
- plA[0] = CurrentOffSetA>=M?0.0:A[0];
- plA[16] = CurrentOffSetA+16>=M?0.0:A[16];
+ plA[0] = CurrentOffSetA>=M?0.0f:A[0];
+ plA[16] = CurrentOffSetA+16>=M?0.0f:A[16];
barrier(CLK_LOCAL_MEM_FENCE);
--- a/src/library/blas/AutoGemm/UserGemmKernelSources/sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_src.cpp
+++ b/src/library/blas/AutoGemm/UserGemmKernelSources/sgemm_Col_TN_B1_MX032_NX032_KX16_BRANCH_src.cpp
@@ -85,11 +85,11 @@ __kernel void sgemm_Col_TN_B1_MX032_NX03
__local float* plB = lB + idx*33+idy;
barrier(CLK_LOCAL_MEM_FENCE);
- plB[0] = CurrentOffSetB>=N?0.0:B[0];
- plB[16] = CurrentOffSetB+16>=N?0.0:B[16*ldb];
+ plB[0] = CurrentOffSetB>=N?0.0f:B[0];
+ plB[16] = CurrentOffSetB+16>=N?0.0f:B[16*ldb];
- plA[0] = CurrentOffSetA>=M?0.0:A[0];
- plA[16] = CurrentOffSetA+16>=M?0.0:A[16*lda];
+ plA[0] = CurrentOffSetA>=M?0.0f:A[0];
+ plA[16] = CurrentOffSetA+16>=M?0.0f:A[16*lda];
barrier(CLK_LOCAL_MEM_FENCE);
--- a/src/library/blas/gens/clTemplates/rotg.cl
+++ b/src/library/blas/gens/clTemplates/rotg.cl
@@ -68,7 +68,7 @@ __kernel void %PREFIXrotg_kernel( __global %TYPE *_A, __global %TYPE *_B, __glob
Creg = Areg / R;
Sreg = Breg / R;
Z = (isgreater(absA, absB))? Sreg:
- ( (isnotequal(Creg, ZERO))? (1.0/Creg): 1.0 );
+ ( (isnotequal(Creg, ZERO))? ((%TYPE)1.0f/Creg): 1.0f );
}
_A[offa] = R;
_B[offb] = Z;
--- a/src/library/blas/gens/clTemplates/rotmg.cl
+++ b/src/library/blas/gens/clTemplates/rotmg.cl
@@ -32,7 +32,7 @@ static const char *rotmg_kernel = "
#define GAM (%TYPE)4096.0
#define GAMSQ (%TYPE)( GAM * GAM )
-#define RGAMSQ (%TYPE)( 1.0 / GAMSQ )
+#define RGAMSQ (%TYPE)( (%TYPE)1.0f / GAMSQ )
__kernel void %PREFIXrotmg_kernel( __global %TYPE *_D1, __global %TYPE *_D2, __global %TYPE *_X1,
__global %TYPE *_Y1, __global %TYPE *_param,
--- a/src/library/blas/gens/clTemplates/syr2_her2.cl
+++ b/src/library/blas/gens/clTemplates/syr2_her2.cl
@@ -172,7 +172,7 @@ __kernel void %PREFIXsyr2_her2_CL_kernel( __global %TYPE* _A, __global const %TY
#ifdef HER2_ONLY
/* HER2 defn: On output, if alpha not equal to 0.0, then imaginary part of A is set to zero. */
- res2.odd = (r == c) ? 0.0 : res2.odd;
+ res2.odd = (r == c) ? 0.0f : res2.odd;
#endif
--- a/src/library/blas/gens/clTemplates/syr_her.cl
+++ b/src/library/blas/gens/clTemplates/syr_her.cl
@@ -45,7 +45,7 @@ __kernel void %PREFIXsyr_her_CL_kernel( __global %TYPE* _A, __global const %TYPE
__local %TYPE yShared[%TARGET_ROWS];
// If alpha is zero, the computation yields to a zero value and therefore doesnot update the A matrix.
- if(alpha == 0.0)
+ if(alpha == 0.0f)
return;
A = _A + offa;
@@ -140,7 +140,7 @@ __kernel void %PREFIXsyr_her_CL_kernel( __global %TYPE* _A, __global const %TYPE
these values. On output, the imaginary parts of diagonal elements are set to zero.
*/
- res.odd = (r == c) ? 0.0 : res.odd;
+ res.odd = (r == c) ? 0.0f : res.odd;
#endif
A( r, c ) = res;
@@ -325,7 +325,7 @@ __kernel void %PREFIXsyr_her_CU_kernel( __global %TYPE* _A, __global const %TYPE
__local %TYPE yShared[%TARGET_ROWS];
// If alpha is zero, the computation yields to a zero value and therefore doesnot update the A matrix.
- if(alpha == 0.0)
+ if(alpha == 0.0f)
return;
A = _A + offa;
@@ -419,7 +419,7 @@ __kernel void %PREFIXsyr_her_CU_kernel( __global %TYPE* _A, __global const %TYPE
these values. On output, the imaginary parts of diagonal elements are set to zero.
*/
- res.odd = (r == c) ? 0.0 : res.odd;
+ res.odd = (r == c) ? 0.0f : res.odd;
#endif
A( r , c ) = res;
