[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