Re: [Gegl-developer] [PATCH] Optimize operation box-blur opencl kernel

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_jia 126 com]
Sent: Thursday, January 23, 2014 9:48 AM
To: Victor Oliveira; tom 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.




Yongjia Zhang 


[Date Prev][Date Next]   [Thread Prev][Thread Next]   [Thread Index] [Date Index] [Author Index]