rotor

Barrier and control flow

Discussion created by rotor on Jun 21, 2010
Latest reply on Jul 12, 2010 by nou

Hi everyone,

I have heard lots and also experienced weird behavior of the program with barrier inside a control flow and since the tech docs does not explain well this so I am seeking for your help. I have several questions on this issue:

1)what the signal that the barrier() statement results to the control unit to commit that it is executed? is that an interrupt signal or program counter...etc?

2)can I you this scenario to avoid the problem:

--the illegal version

if(id=get_local_id(0)%4)

{

loc_data[id]=glob_data[id];

barrier(CLK_LOCAL_MEM_FENCE);

}

do_smt_next++;

-the modified version

 

if(id=get_local_id(0)%4)

{

loc_data[id]=glob_data[id];

}

barrier(CLK_LOCAL_MEM_FENCE);

do_smt_next++;

 

---

so if I move the barrier out of the if flow, will the threads that does not enter the if condition wait @ the barrier until the threads that enter the if condition finishing the task inside the if condition?

 

3) what happen with the barrier in branched condition? like this:

 

if(id=get_local_id(0)%4)

{

loc_data[id]=glob_data[id];

barrier(CLK_LOCAL_MEM_FENCE);

}

else

{

loc_data[id]=glob_data[id]+4;

barrier(CLK_LOCAL_MEM_FENCE);

}



---

since the condition branches out so not all threads will enter either IF or ELSE condition --> what happen with the barrier there?

 

Thanks,

Roto

Outcomes