Local memory same for all workgroup

OpenCL
1

I have a problem with local memory, concretely with save value. This is all kernels in file loaded to program.


__kernel void initData(int bitSample, int transmitorCount ,int bufferLen, __global short *sequence)
{

   __local int constData[2];
   __local short mask[(15 * 1023)];

   constData[0] = bitSample;
   constData[1] = bufferLen;

   for (int i = 0; i < transmitorCount * 1023; i++) {
       mask[i] = sequence[i];
   }

   printf("startPlace: %u 	 Resutl: %i = %i
", transmitorCount, constData[0], constData[1]);
   barrier(CLK_LOCAL_MEM_FENCE);
}

__kernel void correlation(__global short *inBuffer, __global int *outBuffer, __local int *constData)
{

   uint startPlace = get_global_id(0);
   uint sequence = get_global_id(1);

   outBuffer[0] = inBuffer[0];

   printf("startPlace: %i Sequence: %i	 Resutl: %i = %i
", startPlace, sequence, constData[0], constData[1]);
}

Program is working and local memory is saved. Initialization call is:

    size_t globaId[3] = { 1, 0, 0 };
    ret = clEnqueueNDRangeKernel(gpuControlData->cmdQueue, gpuControlData->initKernel, 1, NULL, globaId, NULL, 0, NULL, NULL);

But when I start correlation kernel, saved values is in local only in first start. Second and next starts are wrong data. I try to init local memory before every start, but it doesn´t work. Create command to queue is:

while(true){
    ret = clSetKernelArg(gpuControlData->corelationKernel, 0, sizeof(cl_mem), (void *) &inBuffer);
    ret = clSetKernelArg(gpuControlData->corelationKernel, 1, sizeof(cl_mem), (void *)&outBuffer);
    ret = clSetKernelArg(gpuControlData->corelationKernel, 2, sizeof(cl_int) * 2, NULL);

    size_t globaId[3] = { 3, 2, 0 };
    ret = clEnqueueNDRangeKernel(gpuControlData->maskCorrelQueue, gpuControlData->corelationKernel, 2, NULL, globaId, NULL, 0, NULL, NULL);
}

The result of one run is:


startPlace: 2    Resutl: 100 = 150000
startPlace: 0 Sequence: 0        Resutl: 100 = 150000
startPlace: 1 Sequence: 0        Resutl: 100 = 150000
startPlace: 2 Sequence: 0        Resutl: 100 = 150000
startPlace: 0 Sequence: 1        Resutl: 100 = 150000
startPlace: 1 Sequence: 1        Resutl: 100 = 150000
startPlace: 2 Sequence: 1        Resutl: 100 = 150000
oneIteration
startPlace: 0 Sequence: 0        Resutl: 1056562655 = 1058300571
startPlace: 1 Sequence: 0        Resutl: 1056562655 = 1058300571
startPlace: 2 Sequence: 0        Resutl: 1056562655 = 1058300571
startPlace: 0 Sequence: 1        Resutl: 1056562655 = 1058300571
startPlace: 1 Sequence: 1        Resutl: 1056562655 = 1058300571
startPlace: 2 Sequence: 1        Resutl: 1056562655 = 1058300571

I understand the local memory is saved in only one work group so I need to save that data to every work group.

So I need some help or example how to create same local variable in every work group?
I have 32kB size of lccal memory.

Thanks for help.

2

I think you might be misunderstanding what local memory is for.

Local memory is designed to be used for communication inside of a single work group, during execution of a single kernel. Once a work group finishes execution, the associated local memory is liberated, and attempting to access it in another kernel execution is undefined behaviour (you may encounter the same values as before, or something else entirely if something has overwritten the memory since).

If you want to transmit data between two kernel executions, you will need to use global memory.