Hi yongjia,

Right, the size is indeed uniform for all work items in the same group. I just 
guess people may be a little bit concern of the y which is derived from 
get_local_id(1).
You may consider to add the following attribute to your fast kernel to 
eliminate those concerns.

__attribute__((reqd_work_group_size(256, 1, 1)))


From: 张勇加 [mailto:zhang_yong_...@126.com]
Sent: Thursday, January 23, 2014 9:48 AM
To: Victor Oliveira; t...@stellard.net
Cc: gegl-developer-list@gnome.org; Zou, Nanhai; Gong, Zhigang; Zhang, Yongjia
Subject: Re:Re: [Gegl-developer] [PATCH] Optimize operation box-blur opencl 
kernel

Hi Victor and Tom,

Though the barrier you were talking about is not guaranteed to be executed by 
all threads, but the work items in a workgroup's behaviors about this barrier 
are all the same, so, if a work item executes the barrier, then the others of 
the same work group execute it as well, and vice versa. It's for two reasons. 
1st is the 'size' viarable are all the same value in all work items and it 
decreases without any condition. 2nd is that if the 'break' is called because 
of the  'y==height-1', it is because that we were dealing with the last row of 
the buffer, the work items of a work group are all in the same situation, every 
work item will pass this barrier and return.

Thanks

Yongjia Zhang




_______________________________________________
gegl-developer-list mailing list
List address:    gegl-developer-list@gnome.org
List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list

Reply via email to