Persistent thread on the CPU

Persistent thread on the CPU

Hi,I'm trying to implement some "persistent thread" on the CPU to batch a set of tasks, but I got some strange results.I have put some "printf" in the following code. What is strange is that I see the "BEFORE" before the "START" !I have a local barrier and so I should see "START" before !!

{ const size_t lid = get_local_id(0); __local volatile int localPoolNextRay[1]; __local volatile int localPoolRayCount[1]; if (lid < 1){ localPoolNextRay[0] = localPoolRayCount[0] = 0; printf("START %d %d : %d\\n", get_global_id(0), lid, localPoolNextRay[0]); } barrier(CLK_LOCAL_MEM_FENCE); while(true) { // Local pool is empty if (localPoolRayCount[0] < 1 && lid < 1) { localPoolNextRay[0] = atom_add(globalPoolNextRay, LOAD_BALANCER_BATCH_SIZE); localPoolRayCount[0] = LOAD_BALANCER_BATCH_SIZE; } mem_fence(CLK_LOCAL_MEM_FENCE); //barrier(CLK_LOCAL_MEM_FENCE); printf("BEFORE %d %d : %d\\n", get_global_id(0), lid, localPoolNextRay[0]); // Get rays from local pool int myRayIndex = localPoolNextRay[0] + lid; if (myRayIndex >= globalPoolRayCount) return; printf("AFTER %d\\n", myRayIndex); mem_fence(CLK_LOCAL_MEM_FENCE); if (lid < 1) { localPoolNextRay[0] += 32; localPoolRayCount[0] -= 32; //mem_fence(CLK_LOCAL_MEM_FENCE); } mem_fence(CLK_LOCAL_MEM_FENCE); // Execute trace(myRayIndex, tasks); }}

5 posts / 0 new
Last post
For more complete information about compiler optimizations, see our Optimization Notice.

Hi,

How many workgroups you have?

Can you put small reproducer of this issue?

Thansk,
Evgeny

Hi,Global work = 704Local work = 32I'm able to run it correctly, but for this I have to put barrier everywhere ! In fact, because there is no "natural" SIMT behavior on the CPU I'm not sure that using persitent thread will help on the CPU ! (Lot of barriers or atomic functions !!).

Hi Polar01,

Since you have serveral work-groups executing simultaneously, the 'BEFORE' printout could come from another work group (from a work item that has lid >= 1 and therefore doesn't print 'START')

Good to here it's running well.

Leave a Comment

Please sign in to add a comment. Not a member? Join today