Bug 66931

Summary:	Can't specify local buffer size using clSetKernelArg
Product:	Mesa	Reporter:	Jonathan Charest <jcharest>
Component:	Other	Assignee:	mesa-dev
Status:	RESOLVED FIXED	QA Contact:
Severity:	normal
Priority:	medium
Version:	git
Hardware:	Other
OS:	Linux (All)
Whiteboard:
i915 platform:		i915 features:
Attachments:	Addess spaces support for kernel arguments Addess spaces support for kernel arguments Addess spaces support for kernel arguments

Description Jonathan Charest 2013-07-15 15:22:01 UTC

I always get return code -51 when I try to set the second argument of my kernel. This argument is a local buffer. Here is the call I make:

clSetKernelArg(in_Kernel.mKernel, 1, 1024, NULL);

The kernel prototype is:
__kernel void GetEfficiency( __constant const SimulationParams * const in_pSimulationParams, __local float * l_afEfficiencies, __global const float * const in_pRandomValues, __global float * const out_pfEfficiencies)

So argument index 1 is really a local buffer as can be seen from the prototype. I will try to investigate this further.

Comment 1 Jonathan Charest 2013-07-15 20:17:39 UTC

Created attachment 82462 [details]
Addess spaces support for kernel arguments

Use the data from getAddressSpaceMap to determine the correct type of the kernel arguments.

Comment 2 Tom Stellard 2013-07-16 01:45:44 UTC

(In reply to comment #1)
> Created attachment 82462 [details]
> Addess spaces support for kernel arguments
> 
> Use the data from getAddressSpaceMap to determine the correct type of the
> kernel arguments.

This patch looks good to me, can you send it to mesa-dev@lists.freedesktop.org.
I think for r600g to work correctly, you will also need to make sure we are writing the correct lds size to the SQ_LDS_ALLOC register.  I think what we want is the size reported by clover plus the size reported by the compiler.

Comment 3 Jonathan Charest 2013-07-16 15:15:18 UTC

Created attachment 82485 [details] [review]
Addess spaces support for kernel arguments

Added the local argument size to lds alloc size.

Comment 4 Jonathan Charest 2013-07-16 19:32:57 UTC

Created attachment 82498 [details] [review]
Addess spaces support for kernel arguments

Fixed the local size to be in dw

Comment 5 Tom Stellard 2013-07-19 15:11:13 UTC

(In reply to comment #4)
> Created attachment 82498 [details] [review] [review]
> Addess spaces support for kernel arguments
> 
> Fixed the local size to be in dw

This patch looks good to me, can you send all the paches to mesa-dev@lists.freedesktop.org for review?

Comment 6 Tom Stellard 2014-08-21 15:56:54 UTC

This was fixed in git a while ago.
commit 4f8048bb5a8558ae4313b12ffd70b593cc629fe8
commit d9576598c7e1c6e4fee913a918345190248a9d19

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.