topic Persistent thread on the CPU in OpenCL* for CPU

Persistent thread on the CPU

Polar01 — Mon, 19 Sep 2011 10:59:26 GMT

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);
}
}

Persistent thread on the CPU

Evgeny_F_Intel — Mon, 19 Sep 2011 11:17:15 GMT

Hi,

How many workgroups you have?

Can you put small reproducer of this issue?

Thansk,
Evgeny

Persistent thread on the CPU

Polar01 — Mon, 19 Sep 2011 13:52:41 GMT

Hi,

Global work = 704

Local work = 32

I'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 !!).

Persistent thread on the CPU

Eli_Bendersky__Intel — Tue, 20 Sep 2011 04:53:48 GMT

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')

Persistent thread on the CPU

Evgeny_F_Intel — Tue, 20 Sep 2011 05:43:56 GMT

Good to here it's running well.