smallGame

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,

Hi,

I wanted to generate geometry such as a sphere on the geometry shader, (by sending only a point on the GPU).
But I realized we can only generate 1024 bytes with a geometry shader and that the geometry shader is not designed to generate a lot of data.

That's too bad !! Why a such design ??

Then I thought doesn't matter I am going to use the Hull and Domain shader to tesselate my sphere. But once again I was disppointed the geometry shader is after tesselator stage !!! Again : Why a such design ??

So I guess what I have to do is just generate my sphere on CPU and tesselate it on the tesselator stage; but I d like to understand why they designed the graphic pipeline that way... I am sure I am missing something...

Cheers,