Jump to content

  • Log In with Google      Sign In   
  • Create Account

FREE SOFTWARE GIVEAWAY

We have 4 x Pro Licences (valued at $59 each) for 2d modular animation software Spriter to give away in this Thursday's GDNet Direct email newsletter.


Read more in this forum topic or make sure you're signed up (from the right-hand sidebar on the homepage) and read Thursday's newsletter to get in the running!


OPenCL global memory read write understanding


Old topic!
Guest, the last post of this topic is over 60 days old and at this point you may not reply in this topic. If you wish to continue this conversation start a new topic.

  • You cannot reply to this topic
2 replies to this topic

#1 smallGame   Members   -  Reputation: 208

Like
0Likes
Like

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 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,

Edited by smallGame, 13 December 2012 - 10:28 AM.


Sponsor:

#2 lerlinghagen   Members   -  Reputation: 257

Like
2Likes
Like

Posted 13 December 2012 - 02:35 PM

There is no intra-workgroup synchronization within a kernel. The barrier
instruction only performs synchronization within a given workgroup. The OpenCL specification does not guarantee consistency between different workgroups - see section 3.3.1 of the specification.

#3 smallGame   Members   -  Reputation: 208

Like
0Likes
Like

Posted 13 December 2012 - 04:08 PM

Thanks lerlinghagen Posted Image !!




Old topic!
Guest, the last post of this topic is over 60 days old and at this point you may not reply in this topic. If you wish to continue this conversation start a new topic.



PARTNERS