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@xxxxxxx]
Sent: Thursday, January 23, 2014 9:48 AM
To: Victor Oliveira; tom@xxxxxxxxxxxx
Cc: gegl-developer-list@xxxxxxxxx; Zou, Nanhai; Gong, Zhigang; Zhang, Yongjia
Subject: Re:Re: [PATCH] Optimize operation box-blur opencl kernel
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.
|
_______________________________________________
gegl-developer-list mailing list
List address: gegl-developer-list@xxxxxxxxx
List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list
