[Beignet] [PATCH] Remove global barrier assert.

Zou, Nanhai nanhai.zou at intel.com
Fri Jun 14 00:27:37 PDT 2013


1st,  OCL spec only require order in work item not work group in barrier 

>   dst[get_local_size(0) * (2 * get_group_id(0)) + get_local_id(0)] =
> src[get_local_size(0) * 2 * get_group_id(0) + get_local_size(0) -
(get_local_id(0)
> + 1)];

It seems that you have referenced result from difference thread here.
So you need a barrier between read and write.

Please understand the mem fence in a thread-group semantic in multithread case, 

Say,

Thread A 
Read addr
Write addr

If Thread B see in LLC _Thread_ A write happen before _Thread_ A read.

That is a mem fence issue.

Thanks
Zou Nanhai






More information about the Beignet mailing list