Hello,
I think I have a problem with finding/setting the local buffer but I´m not sure how to describe/explain my problem.
Ok… my situation:
I have for example 10000 Items to be handled and I want to make use of __local memory to speed up. Therefore I create local buffer f.e. like this
__kernel void test(__global MyStruct in_MyStruct)
{
__local MyStruct l_SharedMemory[FIXED_SIZE];
//copy global to local
...
// create barrier
barrier(SOME_FENCE);
// make calculations
..
}
FIXED_SIZE depends on the hardware I´m running. Therefore I thought get CL_KERNEL_WORK_GROUP_SIZE with clGetKernelWorkGroupInfo and use the result to find the optimal value for FIXED_SIZE. Afterwards passing FIXED_SIZE as a compiler option when calling clCreateKernel. Theoretically a good idea - the problem is that for clGetKernelWorkGroupInfo I need a kernel which isn´t created at that time. Any ideas how to find a correct value for fixed size?!?!?!
You are right, when the runtime determines preferred workgroup size of the kernel, it only looks at register usage. Local memory allocation inside the kernel can only be taken into account, if the size is a compile-time constant (let’s say a #define SIZE 128 that is written at the beginning of the code by the host application). However, my findings are that local memory allocation is not taken into account when determining preferred WGS.
Easiest solution is to query local memory size by clGetDeviceInfo(), ask the kernel for preferred WGS based on register usage, and allocate local memory inside host code by taking the greatest common factor, so to say. With these two informations you can find optimal parameters for your application.
Only trick (or convenience) is to allocate local memory by the host, since that is where all the neccesary information is gathered in the first place.
Ok… the samples mostly use a hard coded value - so it wasn´t obvious what you were trying say me. But after inverstigating a little bit… I got an idea. Thanks for the hint… I think it´s not the best and cleanest solution but I will give it try. Thnaks a lot…