[Beignet] [PATCH] Remove global barrier assert.

[Zhigang Gong]

> -----Original Message-----
> From: beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org
>
[mailto:beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org]
> On Behalf Of Zou, Nanhai
> Sent: Friday, June 14, 2013 3:28 PM
> To: Zhigang Gong; Yang, Rong R
> Cc: beignet at lists.freedesktop.org
> Subject: Re: [Beignet] [PATCH] Remove global barrier assert.
> 
> 1st,  OCL spec only require order in work item not work group in barrier
[Gong, Zhigang] This is the major difference between us. I insist that a
barrier
with global memory fence need to make sure all the instructions before the
barrier have been executed, and all the global memory operations before the
barrier have been processed. And then it can start to execute the following
instructions. And that can make sure we can synchronize different work items
within one work group to share data. Otherwise, there will be no way to
share
data between work items in one work group. Right?

> 
> >   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.
[Gong, Zhigang] There is a barrier with global memory fence flag.

> 
> 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
> 
> 
> 
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet