Bug#921890: marked as done (clblas: error: variables in the local address space can only be declared in the outermost scope of a kernel function)

Thu, 28 Sep 2023 18:09:15 -0700 

Your message dated Fri, 29 Sep 2023 01:04:00 +0000
with message-id <[email protected]>
and subject line Bug#921890: fixed in clblas 2.12-4
has caused the Debian Bug report #921890,
regarding clblas: error: variables in the local address space can only be 
declared in the outermost scope of a kernel function
to be marked as done.

This means that you claim that the problem has been dealt with.
If this is not the case it is now your responsibility to reopen the
Bug report if necessary, and/or fix the problem forthwith.

(NB: If you are a system administrator and have no idea what this
message is talking about, this may indicate a serious mail system
misconfiguration somewhere. Please contact [email protected]
immediately.)


-- 
921890: https://bugs.debian.org/cgi-bin/bugreport.cgi?bug=921890
Debian Bug Tracking System
Contact [email protected] with problems
--- Begin Message --- 
Package: clblas
Version: 2.12-1
Control: tags -1 patch upstream
(upstream tag based on checking the source, not testing)

Some clblas operations fail on beignet-opencl-icd with
stringInput.cl:179:24: error: variables in the local address space can only be declared in the outermost scope of a kernel function
Checking the source confirms that they are declaring __local variables in an inner scope, which is not allowed by the OpenCL standard: https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/local.html
This affects example_chpmv, example_snrm2, example_sspmv, example_stpmv, example_stpsv, example_strmv, example_strsv.
The attached patch makes the examples run, but has *not* been tested beyond them (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: Move __local declarations to kernel function scope

The OpenCL spec does not allow declaring local variables in scopes
below kernel function scope, and such declarations fail to build
on at least beignet-opencl-icd.

https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/local.html

Author: Rebecca N. Palmer <[email protected]>
Bug-Debian: https://bugs.debian.org/<bugnumber>
Forwarded: no

--- clblas-2.12.orig/src/library/blas/gens/clTemplates/trmv.cl
+++ clblas-2.12/src/library/blas/gens/clTemplates/trmv.cl
@@ -75,6 +75,8 @@ __kernel void %PREFIXtrmv_CU_kernel( __g
 
 
 	__local %TYPE  sXData[ TARGET_WIDTH ]; // Each column is multiplied with a common x_vector element
+    volatile __local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * TARGET_WIDTH];
+    volatile __local %TYPE* sData = sDataTemp;
 
 	const int gIdx = get_global_id(0);
 	const int bIdx = get_group_id(0);
@@ -197,8 +199,6 @@ __kernel void %PREFIXtrmv_CU_kernel( __g
 		}
 
 
-		volatile __local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * TARGET_WIDTH];
-		volatile __local %TYPE* sData = sDataTemp;
 		//sDataTemp[(threadIdx & ( TARGET_ROWS_BY_VEC -1 )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		sDataTemp[(threadIdx % ( TARGET_ROWS_BY_VEC )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		barrier(CLK_LOCAL_MEM_FENCE);
@@ -325,6 +325,8 @@ __kernel void %PREFIXtrmv_CL_kernel( __g
     #endif
 
 	__local %TYPE sXData[ TARGET_WIDTH ]; // Each column is multiplied with a common x_vector element
+    __local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * TARGET_WIDTH];
+    __local %TYPE* sData = sDataTemp;
 
 	size_t gIdx = get_global_id(0);
 	size_t bIdx = get_group_id(0);
@@ -448,8 +450,6 @@ __kernel void %PREFIXtrmv_CL_kernel( __g
 		}
 
 
-		__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * TARGET_WIDTH];
-		__local %TYPE* sData = sDataTemp;
 		//sDataTemp[(threadIdx & ( TARGET_ROWS_BY_VEC -1 )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		sDataTemp[(threadIdx % ( TARGET_ROWS_BY_VEC  )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		barrier(CLK_LOCAL_MEM_FENCE);
@@ -584,6 +584,7 @@ __kernel void %PREFIXtrmv_CLT_kernel( __
 	int threadIdx 	= get_local_id(0);
 
 	__local %TYPE xShared[TARGET_WIDTH];
+    __local %TYPE%V* xSharedTemp;
 
 	int startCol  	= blockIdx * %TARGET_ROWS;
 
@@ -608,7 +609,7 @@ __kernel void %PREFIXtrmv_CLT_kernel( __
 
 		//float4 xData = (float4)(xShared[ rowShift ], xShared[ rowShift + 1], xShared[ rowShift + 2], xShared[ rowShift + 3]);
 		%TYPE%V xData;
-		__local %TYPE%V* xSharedTemp = (xShared + rowShift);
+		xSharedTemp = (xShared + rowShift);
 		xData = *(xSharedTemp);
 
 		int row	= startRow + rowShift;
@@ -761,6 +762,9 @@ __kernel void %PREFIXtrmv_CUT_kernel( __
 	int threadIdx 	= get_local_id(0);
 
 	__local %TYPE xShared[TARGET_WIDTH];
+    __local %TYPE%V* xSharedTemp;
+    __local %TYPE%V sDataTemp[TARGET_WIDTH_BY_VEC * %TARGET_ROWS];
+    __local %TYPE* sData = sDataTemp;
 
 	int startRow  	= 0;
 	int startCol  	= N - (blockIdx + 1)* %TARGET_ROWS;
@@ -842,7 +846,7 @@ __kernel void %PREFIXtrmv_CUT_kernel( __
 
 			//float4 xData = (float4)(xShared[ rowShift ], xShared[ rowShift + 1], xShared[ rowShift + 2], xShared[ rowShift + 3]);
 			%TYPE%V xData;
-			__local %TYPE%V* xSharedTemp = (xShared + rowShift);
+			xSharedTemp = (xShared + rowShift);
 			xData = *(xSharedTemp);
 
 			int row	= startRow + rowShift;
@@ -861,8 +865,6 @@ __kernel void %PREFIXtrmv_CUT_kernel( __
 		//__local float4 sData[16][4];
 		//sData[(threadIdx & 15)][colShift] = acc;
 		//barrier(CLK_LOCAL_MEM_FENCE);
-		__local %TYPE%V sDataTemp[TARGET_WIDTH_BY_VEC * %TARGET_ROWS];
-		__local %TYPE* sData = sDataTemp;
 
 		//sDataTemp[ ( threadIdx & ( TARGET_WIDTH_BY_VEC -1 ) ) + (colShift * TARGET_WIDTH_BY_VEC) ] = acc;
 		sDataTemp[ ( threadIdx % ( TARGET_WIDTH_BY_VEC  ) ) + (colShift * TARGET_WIDTH_BY_VEC) ] = acc;
--- clblas-2.12.orig/src/library/blas/gens/clTemplates/trsv_gemv.cl
+++ clblas-2.12/src/library/blas/gens/clTemplates/trsv_gemv.cl
@@ -35,6 +35,11 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 {
 	__global %TYPE* xnew;
 	__global %TYPE* A = _A + offa;
+	__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
+	//__local %TYPE* sData = sDataTemp;
+	__local %TYPE  xShared_scalar; // To share solved x value with other threads..
+	__local %TYPE xShared_array[%V];
+
 
 	if ( incx < 0 ) // Goto end of vector
 	{
@@ -100,7 +105,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 			// As the above condition ( targetRow <= lastRow) changes targetCol for only threads with condition true
 			targetCol 	= startCol - %TARGET_ROWS;
 
-			__local %TYPE  xShared; // To share solved x value with other threads..
 
 			for( int i=0; i < (lastRow + 1); i++)
 			{
@@ -108,8 +112,8 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 				{
 					%TYPE xVal = xnew[ targetRow * incx];
 					%SUB(sum, xVal, sum);
-					xShared = sum;
-					xnew[ targetRow * incx ] = xShared;
+					xShared_scalar = sum;
+					xnew[ targetRow * incx ] = xShared_scalar;
 				}
 
 				barrier(CLK_LOCAL_MEM_FENCE);
@@ -118,7 +122,7 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 				{
 					loadedA = A((targetRow), (targetCol));
 					%CONJUGATE(doConj, loadedA);
-					%MAD(sum, loadedA, xShared);
+					%MAD(sum, loadedA, xShared_scalar);
 				}
 
 				// Avoid Race
@@ -164,8 +168,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 			barrier(CLK_LOCAL_MEM_FENCE);
 		}
 
-		__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
-		//__local %TYPE* sData = sDataTemp;
 		sDataTemp[(threadIdx % ( TARGET_ROWS_BY_VEC )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		barrier(CLK_LOCAL_MEM_FENCE);
 
@@ -193,7 +195,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 				}
 			}
 
-			__local %TYPE xShared[%V];
 
 			int targetRowTemp = rowStart + threadIdx * %V;
 			int VECTOR_SIZE   = %V;
@@ -231,7 +232,7 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 					// Solve for first x - Do the rest in loop
 					%TYPE x[%V];
 					%SUB(x[VECTOR_SIZE - 1], xVal[VECTOR_SIZE - 1], sumVecReg[VECTOR_SIZE - 1]);
-					xShared[%V - 1] = x[%V - 1];
+					xShared_array[%V - 1] = x[%V - 1];
 					xnew[ (targetRowTemp + %V - 1)* incx ] = x[%V - 1];
 
 					//#pragma unroll
@@ -256,13 +257,13 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 					//#pragma unroll
 					for(int m = 0; m < %V; m++)
 					{
-						xShared[m] = x[m];
+						xShared_array[m] = x[m];
 						xnew[ (targetRowTemp + m)* incx ] = x[m];
 					}
 			    }
 
 
-			    // Sync so that xShared it available to all threads
+			    // Sync so that xShared_array is available to all threads
 			    barrier(CLK_LOCAL_MEM_FENCE);
 
 			    if ( threadIdx < (TARGET_ROWS_BY_VEC - 1 - i))
@@ -270,10 +271,10 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 						//#pragma unroll
 						for( int j=0; j < %V; j++)
 						{
-							//sumVec += vload4( 0, &A((targetRowTemp), (targetCol -j))) * xShared[%V - 1 -j];
+							//sumVec += vload4( 0, &A((targetRowTemp), (targetCol -j))) * xShared_array[%V - 1 -j];
 							%TYPE%V loadedAVec  = %VLOAD( 0, &A((targetRowTemp), (targetCol -j)));
 							%CONJUGATE(doConj, loadedAVec);
-							%VMAD(sumVec, loadedAVec, xShared[VECTOR_SIZE - 1 -j]);
+							%VMAD(sumVec, loadedAVec, xShared_array[VECTOR_SIZE - 1 -j]);
 						}
 				}
 
@@ -335,6 +336,11 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 {
 	__global %TYPE* xnew;
 	__global %TYPE* A = _A + offa;
+	__local %TYPE xData[ %TARGET_WIDTH];
+	__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
+	//__local %TYPE* sData = sDataTemp;
+	__local %TYPE  xShared_scalar; // To share solved x value with other threads..
+	__local %TYPE xShared_array[%V];
 
 	if ( incx < 0 ) // Goto end of vector
 	{
@@ -400,7 +406,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 			// As the above condition ( targetRow <= lastRow) changes targetCol for only threads with condition true
 			targetCol 	= startCol - %TARGET_ROWS;
 
-			__local %TYPE  xShared; // To share solved x value with other threads..
 
 			for( int i=0; i < (lastRow + 1); i++)
 			{
@@ -412,9 +417,9 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 					// Handle diagonal element
 					loadedA = A((targetRow), (targetCol));
 					%CONJUGATE(doConj, loadedA);
-					%DIV(xShared, sum, loadedA);
+					%DIV(xShared_scalar, sum, loadedA);
 
-					xnew[ targetRow * incx ] = xShared;
+					xnew[ targetRow * incx ] = xShared_scalar;
 				}
 
 				barrier(CLK_LOCAL_MEM_FENCE);
@@ -423,7 +428,7 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 				{
 					loadedA = A((targetRow), (targetCol));
 					%CONJUGATE(doConj, loadedA);
-					%MAD(sum, loadedA, xShared);
+					%MAD(sum, loadedA, xShared_scalar);
 				}
 
 				// Avoid Race
@@ -443,7 +448,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 		%TYPE   sumTemp = %MAKEVEC(0.0);
 		%TYPE%V sum	= %VMAKEVEC(sumTemp);
 
-		__local %TYPE xData[ %TARGET_WIDTH];
 
 		//#pragma unroll
 		for( int i=1; i <= %NLOOPS; i++)
@@ -469,8 +473,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 			barrier(CLK_LOCAL_MEM_FENCE);
 		}
 
-		__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
-		//__local %TYPE* sData = sDataTemp;
 		sDataTemp[(threadIdx % ( TARGET_ROWS_BY_VEC )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		barrier(CLK_LOCAL_MEM_FENCE);
 
@@ -498,7 +500,6 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 				}
 			}
 
-			__local %TYPE xShared[%V];
 
 			int targetRowTemp = rowStart + threadIdx * %V;
 			int VECTOR_SIZE   = %V;
@@ -538,7 +539,7 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 					%SUB(x[VECTOR_SIZE - 1], xVal[VECTOR_SIZE - 1], sumVecReg[VECTOR_SIZE - 1]);
 					%DIV(sumVecReg[VECTOR_SIZE - 1], x[VECTOR_SIZE -1], reg[VECTOR_SIZE - 1][VECTOR_SIZE - 1]);
 					x[VECTOR_SIZE -1] = sumVecReg[VECTOR_SIZE - 1];
-					xShared[%V - 1] = x[%V - 1];
+					xShared_array[%V - 1] = x[%V - 1];
 					xnew[ (targetRowTemp + %V - 1)* incx ] = x[%V - 1];
 
 					//#pragma unroll
@@ -568,13 +569,13 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 					//#pragma unroll
 					for(int m = 0; m < %V; m++)
 					{
-						xShared[m] = x[m];
+						xShared_array[m] = x[m];
 						xnew[ (targetRowTemp + m)* incx ] = x[m];
 					}
 			        }
 
 
-			        // Sync so that xShared it available to all threads
+			        // Sync so that xShared_array is available to all threads
 			        barrier(CLK_LOCAL_MEM_FENCE);
 
 			      	if ( threadIdx < (TARGET_ROWS_BY_VEC - 1 - i))
@@ -582,10 +583,10 @@ __kernel void %PREFIXtrsv_CU_ComputeRect
 						//#pragma unroll
 						for( int j=0; j < %V; j++)
 						{
-							//sumVec += vload4( 0, &A((targetRowTemp), (targetCol -j))) * xShared[%V - 1 -j];
+							//sumVec += vload4( 0, &A((targetRowTemp), (targetCol -j))) * xShared_array[%V - 1 -j];
 							%TYPE%V loadedAVec  = %VLOAD( 0, &A((targetRowTemp), (targetCol -j)));
 							%CONJUGATE(doConj, loadedAVec);
-							%VMAD(sumVec, loadedAVec, xShared[VECTOR_SIZE - 1 -j]);
+							%VMAD(sumVec, loadedAVec, xShared_array[VECTOR_SIZE - 1 -j]);
 						}
 				}
 
@@ -649,6 +650,11 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 {
 	__global %TYPE* xnew;
 	__global %TYPE* A = _A + offa;
+	__local %TYPE xData[ %TARGET_WIDTH];
+	__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
+	//__local %TYPE* sData = sDataTemp;
+	__local %TYPE  xShared_scalar; // To share solved x value with other threads..
+	__local %TYPE xShared_array[%V];
 
 	if ( incx < 0 ) // Goto end of vector
 	{
@@ -715,7 +721,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 			// As the above condition ( targetRow <= lastRow) changes targetCol for only threads with condition true
 			targetCol 	= startCol + %TARGET_ROWS;
 
-			__local %TYPE  xShared; // To share solved x value with other threads..
 
 			for( int i=0; i < ((lastRow -startRow) + 1); i++)
 			{
@@ -726,16 +731,16 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 
 					if( isUnity)
 					{
-						xShared = sum;
+						xShared_scalar = sum;
 					}
 					else // Handle diagonal element
 					{
 						loadedA = A((targetRow), (targetCol));
 						%CONJUGATE(doConj, loadedA);
-						%DIV(xShared, sum, loadedA);
+						%DIV(xShared_scalar, sum, loadedA);
 					}
 
-					xnew[ targetRow * incx ] = xShared;
+					xnew[ targetRow * incx ] = xShared_scalar;
 				}
 
 				barrier(CLK_LOCAL_MEM_FENCE);
@@ -744,7 +749,7 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 				{
 					loadedA = A((targetRow), (targetCol));
 					%CONJUGATE(doConj, loadedA);
-					%MAD(sum, loadedA, xShared);
+					%MAD(sum, loadedA, xShared_scalar);
 				}
 
 				// Avoid Race
@@ -764,7 +769,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 		%TYPE   sumTemp = %MAKEVEC(0.0);
 		%TYPE%V sum	= %VMAKEVEC(sumTemp);
 
-		__local %TYPE xData[ %TARGET_WIDTH];
 
 		//#pragma unroll
 		for( int i=1; i <= %NLOOPS; i++)
@@ -790,8 +794,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 			barrier(CLK_LOCAL_MEM_FENCE);
 		}
 
-		__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
-		//__local %TYPE* sData = sDataTemp;
 		sDataTemp[(threadIdx % ( TARGET_ROWS_BY_VEC )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		barrier(CLK_LOCAL_MEM_FENCE);
 
@@ -819,7 +821,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 				}
 			}
 
-			__local %TYPE xShared[%V];
 
 			int targetRowTemp = rowStart + threadIdx * %V;
 			int VECTOR_SIZE   = %V;
@@ -857,7 +858,7 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 					// Solve for first x - Do the rest in loop
 					%TYPE x[%V];
 					%SUB(x[0], xVal[0], sumVecReg[0]);
-					xShared[0] = x[0];
+					xShared_array[0] = x[0];
 					xnew[ (targetRowTemp)* incx ] = x[0];
 
 					//#pragma unroll
@@ -882,13 +883,13 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 					//#pragma unroll
 					for(int m = 0; m < %V; m++)
 					{
-						xShared[m] = x[m];
+						xShared_array[m] = x[m];
 						xnew[ (targetRowTemp + m)* incx ] = x[m];
 					}
 			        }
 
 
-			        // Sync so that xShared it available to all threads
+			        // Sync so that xShared_array is available to all threads
 			        barrier(CLK_LOCAL_MEM_FENCE);
 			      	if ( (threadIdx > i) && ( threadIdx < (TARGET_ROWS_BY_VEC)) )
 				{
@@ -897,7 +898,7 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 						{
 							%TYPE%V loadedAVec  = %VLOAD( 0, &A((targetRowTemp), (targetCol +j)));
 							%CONJUGATE(doConj, loadedAVec);
-							%VMAD(sumVec, loadedAVec, xShared[j]);
+							%VMAD(sumVec, loadedAVec, xShared_array[j]);
 						}
 				}
 
@@ -958,6 +959,12 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 {
 	__global %TYPE* xnew;
 	__global %TYPE* A = _A + offa;
+	__local %TYPE xData[ %TARGET_WIDTH];
+	__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
+	//__local %TYPE* sData = sDataTemp;
+	__local %TYPE  xShared_scalar; // To share solved x value with other threads..
+	__local %TYPE xShared_array[%V];
+
 
 	if ( incx < 0 ) // Goto end of vector
 	{
@@ -1024,7 +1031,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 			// As the above condition ( targetRow <= lastRow) changes targetCol for only threads with condition true
 			targetCol 	= startCol + %TARGET_ROWS;
 
-			__local %TYPE  xShared; // To share solved x value with other threads..
 
 			for( int i=0; i < ((lastRow -startRow) + 1); i++)
 			{
@@ -1036,8 +1042,8 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 					// Handle diagonal element
 					loadedA = A((targetRow), (targetCol));
 					%CONJUGATE(doConj, loadedA);
-					%DIV(xShared, sum, loadedA);
-					xnew[ targetRow * incx ] = xShared;
+					%DIV(xShared_scalar, sum, loadedA);
+					xnew[ targetRow * incx ] = xShared_scalar;
 				}
 
 				barrier(CLK_LOCAL_MEM_FENCE);
@@ -1046,7 +1052,7 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 				{
 					loadedA = A((targetRow), (targetCol));
 					%CONJUGATE(doConj, loadedA);
-					%MAD(sum, loadedA, xShared);
+					%MAD(sum, loadedA, xShared_scalar);
 				}
 
 				// Avoid Race
@@ -1066,7 +1072,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 		%TYPE   sumTemp = %MAKEVEC(0.0);
 		%TYPE%V sum	= %VMAKEVEC(sumTemp);
 
-		__local %TYPE xData[ %TARGET_WIDTH];
 
 		//#pragma unroll
 		for( int i=1; i <= %NLOOPS; i++)
@@ -1092,8 +1097,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 			barrier(CLK_LOCAL_MEM_FENCE);
 		}
 
-		__local %TYPE%V sDataTemp[TARGET_ROWS_BY_VEC * %TARGET_WIDTH];
-		//__local %TYPE* sData = sDataTemp;
 		sDataTemp[(threadIdx % ( TARGET_ROWS_BY_VEC )) + (colShift * TARGET_ROWS_BY_VEC)] = sum;
 		barrier(CLK_LOCAL_MEM_FENCE);
 
@@ -1121,7 +1124,6 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 				}
 			}
 
-			__local %TYPE xShared[%V];
 
 			int targetRowTemp = rowStart + threadIdx * %V;
 			int VECTOR_SIZE   = %V;
@@ -1161,7 +1163,7 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 					%SUB(x[0], xVal[0], sumVecReg[0]);
 					%DIV(sumVecReg[0], x[0], reg[0][0]);
 					x[0] = sumVecReg[0];
-					xShared[0] = sumVecReg[0];
+					xShared_array[0] = sumVecReg[0];
 					xnew[ (targetRowTemp)* incx ] = sumVecReg[0];
 
 					//#pragma unroll
@@ -1191,13 +1193,13 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 					//#pragma unroll
 					for(int m = 1; m < %V; m++)
 					{
-						xShared[m] = x[m];
+						xShared_array[m] = x[m];
 						xnew[ (targetRowTemp + m)* incx ] = x[m];
 					}
 			        }
 
 
-			        // Sync so that xShared it available to all threads
+			        // Sync so that xShared_array is available to all threads
 			        barrier(CLK_LOCAL_MEM_FENCE);
 			      	if ( (threadIdx > i) && ( threadIdx < (TARGET_ROWS_BY_VEC)) )
 				{
@@ -1206,7 +1208,7 @@ __kernel void %PREFIXtrsv_CL_ComputeRect
 						{
 							%TYPE%V loadedAVec  = %VLOAD( 0, &A((targetRowTemp), (targetCol +j)));
 							%CONJUGATE(doConj, loadedAVec);
-							%VMAD(sumVec, loadedAVec, xShared[j]);
+							%VMAD(sumVec, loadedAVec, xShared_array[j]);
 						}
 				}
 
--- clblas-2.12.orig/src/library/blas/gens/clTemplates/nrm2.cl
+++ clblas-2.12/src/library/blas/gens/clTemplates/nrm2.cl
@@ -172,8 +172,14 @@ __kernel void %PREFIXnrm2_ssq_kernel( __
 
     // If scaleOfWG was zero, that means the whole array encountered before was filled with zeroes
     // Note: _scale is a local variable, either all enter or none
-    if(isnotequal(scaleOfWG, PZERO))
+    if(isequal(scaleOfWG, PZERO))
     {
+        if( (get_local_id(0)) == 0 ) {
+            scratchBuff[ get_group_id(0) ] = scaleOfWG;
+            scratchBuff[ numWGs + get_group_id(0) ] = 0.0f;
+        }
+    return;
+    }
         for( gOffset=(get_global_id(0) * %V); (gOffset + %V - 1)<N; gOffset+=( get_global_size(0) * %V ) )
         {
             %TYPE%V vReg1;
@@ -201,7 +207,7 @@ __kernel void %PREFIXnrm2_ssq_kernel( __
         }
 
         %REDUCTION_BY_SUM( ssq );
-    }
+
 
     if( (get_local_id(0)) == 0 ) {
         scratchBuff[ get_group_id(0) ] = scaleOfWG;
--- clblas-2.12.orig/src/library/blas/gens/clTemplates/reduction.cl
+++ clblas-2.12/src/library/blas/gens/clTemplates/reduction.cl
@@ -319,8 +319,13 @@ __kernel void %PREFIXred_ssq_kernel( __g
 
     // If scale was zero, that means the whole array encountered before was filled with zeroes
     // Note: scale is a local variable, either all enter or none
-    if(isnotequal(scaleOfWG, ZERO))
+    if(isequal(scaleOfWG, ZERO))
     {
+        if( (get_local_id(0)) == 0 ) {
+            res[0] = 0.0f;
+        }
+    return;
+    }
         for( gOffset=(get_global_id(0) * %V); (gOffset + %V - 1)<N; gOffset+=( get_global_size(0) * %V ) )
         {
             %TYPE%V scale1, ssq1;
@@ -342,7 +347,7 @@ __kernel void %PREFIXred_ssq_kernel( __g
         }
 
         %REDUCTION_BY_SUM( ssq );
-    }
+
 
     if( (get_local_id(0)) == 0 ) {
         res[0] = scaleOfWG * sqrt(ssq);

--- End Message ---
--- Begin Message --- 
Source: clblas
Source-Version: 2.12-4
Done: tony mancill <[email protected]>

We believe that the bug you reported is fixed in the latest version of
clblas, which is due to be installed in the Debian FTP archive.

A summary of the changes between this version and the previous one is
attached.

Thank you for reporting the bug, which will now be closed.  If you
have further comments please address them to [email protected],
and the maintainer will reopen the bug report if appropriate.

Debian distribution maintenance software
pp.
tony mancill <[email protected]> (supplier of updated clblas package)

(This message was generated automatically at their request; if you
believe that there is a problem with it please contact the archive
administrators by mailing [email protected])


-----BEGIN PGP SIGNED MESSAGE-----
Hash: SHA512

Format: 1.8
Date: Thu, 28 Sep 2023 16:21:54 -0700
Source: clblas
Architecture: source
Version: 2.12-4
Distribution: unstable
Urgency: medium
Maintainer: Debian Science Maintainers 
<[email protected]>
Changed-By: tony mancill <[email protected]>
Closes: 877316 921611 921890
Changes:
 clblas (2.12-4) unstable; urgency=medium
 .
   * Team upload.
 .
   [ Jonathan Bergh ]
   * Fix double literals used in single precision functions (Closes: #877316)
   * Fix crash due to null pointer after compile failure (Closes: #921611)
   * Fix issue where variables are not declared in outermost scope
     (Closes: #921890)
 .
   [ tony mancill ]
   * Use debhelper-compat 13
   * Set Rules-Requires-Root: no in debian/control
   * libclblas-doc now installs docs in /usr/share/doc/libclblas-dev/html
     This seems to be part of the debhelper compat change from 10 to 13
     and matches the behavior of clfft.
   * Update doc-base for actual install directory (libcblas-dev)
   * Document uninstalled files in debian/not-installed
   * Update autopkgtest to use AUTOPKGTEST_TMP
   * Bump Standards-Version to 4.6.2
   * Override dh_auto_install-indep for arch-all (only) build
   * Update autopkgtest for new location of example files
Checksums-Sha1:
 d7e51cf0a984b368a4f3dd566d165d8f3d5fcdfd 2370 clblas_2.12-4.dsc
 ac56ae24186da8aaa96633161edd5253e6bc69cf 12976 clblas_2.12-4.debian.tar.xz
 1eb00f44a90724f83b32a5d66ef2df41d967d268 9301 clblas_2.12-4_amd64.buildinfo
Checksums-Sha256:
 37b5ead62a0859e5a74a73b0fcb730b1daa691051fa10be091bc10b8de38c81d 2370 
clblas_2.12-4.dsc
 e307e73584a1ba7f1140f5fbd524d89aa53697137e012ba8318cdc9b97fb2aa0 12976 
clblas_2.12-4.debian.tar.xz
 98004b5c1108a03e457a9a8101ada610f739cbb58b67083af00937eaa501ebfd 9301 
clblas_2.12-4_amd64.buildinfo
Files:
 adf91cdd262541fc9223e263059b9f40 2370 science optional clblas_2.12-4.dsc
 175f72ec94e70ed5a2231195851d9059 12976 science optional 
clblas_2.12-4.debian.tar.xz
 bfe1927d68f4bcc0322c383c12f1c84e 9301 science optional 
clblas_2.12-4_amd64.buildinfo

-----BEGIN PGP SIGNATURE-----

iQJIBAEBCgAyFiEE5Qr9Va3SequXFjqLIdIFiZdLPpYFAmUWHzMUHHRtYW5jaWxs
QGRlYmlhbi5vcmcACgkQIdIFiZdLPpbuPQ//UNJ0+kkZJ4Xa8vZI5i7XIbwmmKfk
SQEqYiV43p9yd7gi504NBoWVB1UwW+tGClnGOPhGa4dhgulmb5UV1sbDiJUOKrXC
g/Y6Rii0BbDiZ2g8THrzCdJXpJDR5FftZpGpd+ennCgpKDsoccltjEKgqPxE8nZd
HQs/R12+K6VATbKQiskLBlAFn3DUdLyj5AQrXQey6iMrcDTWvtRnz+jcKlvJ4vqt
3pxR8ctulh4FFscsrEupps7laPTVakL+na9pxltgeowxTgHfqfG5TVg2gtCg0Zb0
9D67hGKfs59AkoAFnrFIqcq6rK0pvhUixZqilOkx1c53VMln7qRid7NXK8EzL8b6
v8luuXG5kxiGaNPg4zyDSA0wcbDuu57l8T8BM9cQ++lLz6EHdzqGcu7jWl/beT4q
jVY6vIPjf3p8niDApFTjg5gSKKF6IMLX7r4HOzOPsKK2k9mtAt+CioysCqwt+M1H
PTmJ981E9YhqxG68T+dlvFXiJx33fF9VE/Bmq4fMQONzMd/n2ox/Z/tF++KhNrbc
ZHj3UVQKTfpf10B8QmCLcdvmSjFO/tpBnoa6/54Kq9dlUm+J2mxVRdUbkem/RBGy
gQ6yAbvpA1ha2DU8jQpeduUO60sYToVJRqTRA1Wy2dWhFsMxXWefy2WM23DXipHs
/ZSjZ8X+8wBr/D4=
=1uXf
-----END PGP SIGNATURE-----

--- End Message ---
-- 
debian-science-maintainers mailing list
[email protected]
https://alioth-lists.debian.net/cgi-bin/mailman/listinfo/debian-science-maintainers

Reply via email to