Jump to content

  • Log In with Google      Sign In   
  • Create Account


#ActualsmallGame

Posted 13 December 2012 - 10:28 AM

Hi,

I wrote a simple kernel, I read and write in global memory, but I have some troubles.
From my understanding any work item A could write in the global memory, and any work item B could read what was written by A even if they are not in the same work group. (of course if synchronization is correct)


__kernel void Test(volatile __global int* volatileGlobalMemory) // CL_MEM_READ_WRITE memory flag, size 10
{
	size_t globalId = get_global_id(0);
	// Init array to zero
	if (globalId < 10)
	{
		volatileGlobalMemory[globalId] = 0;
	}
	barrier(CLK_GLOBAL_MEM_FENCE);
	// atomic histogram
	int indexBucket = globalId % 10;
	atomic_inc(&volatileGlobalMemory[indexBucket]);
	// Here everything is fine we have the expected hitogram inside all work items and work group
	barrier(CLK_GLOBAL_MEM_FENCE);
	// Copy just once; this seems to work only on the first work group, I  don't understand why ???!!!
	if (globalId == 0)
	{
		for (int i = 1; i < 10; i++)
		{
			volatileGlobalMemory[i] += volatileGlobalMemory[i - 1];
		}
	}
	barrier(CLK_GLOBAL_MEM_FENCE);
	// This value is correct on the first work group, wrong on the other ones, why ?
	int a0 = volatileGlobalMemory[0];
}

Thanks for your help,

#1smallGame

Posted 13 December 2012 - 10:27 AM

Hi,

I wrote a simple kernel, I read and write in global memory, but I have some troubles.
From my understanding any work item A could write in the global memory, and any work item B could read what was writtenby A even if they are not in the same work group. (of course if synchronization is correct)


__kernel void Test(volatile __global int* volatileGlobalMemory) // CL_MEM_READ_WRITE memory flag, size 10
{
    size_t globalId = get_global_id(0);
    // Init array to zero
    if (globalId < 10)
    {
	    volatileGlobalMemory[globalId] = 0;
    }
    barrier(CLK_GLOBAL_MEM_FENCE);
    // atomic histogram
    int indexBucket = globalId % 10;
    atomic_inc(&volatileGlobalMemory[indexBucket]);
    // Here everything is fine we have the expected hitogram inside all work items and work group
    barrier(CLK_GLOBAL_MEM_FENCE);
    // Copy just once; this seems to work only on the first work group, I  don't understand why ???!!!
    if (globalId == 0)
    {
	    for (int i = 1; i < 10; i++)
	    {
		    volatileGlobalMemory[i] += volatileGlobalMemory[i - 1];
	    }
    }
    barrier(CLK_GLOBAL_MEM_FENCE);
    // This value is correct on the first work group, wrong on the other ones, why ?
    int a0 = volatileGlobalMemory[0];
}

Thanks for your help,

PARTNERS