如何解决write_image是原子的吗?使用atomic_max更好吗?
全面披露:我目前在kronos opencl论坛上发帖,因为到目前为止我还没有收到任何回复:
https://community.khronos.org/t/is-write-image-atomic-is-it-better-than-atomic-max/106418
我正在为图像(2d和3d)编写连接的组件标记算法;我发现没有现有的实现,因此决定基于指针跳转和“回收步骤”来写一个(顺便说一句:如果您知道一个易于使用的,生产就绪的连接组件标签,请告诉我)。
用于2d图像的“ recollection”步骤内核伪代码如下:
1) global_id = (x,y)
2) read v from img[x,y],decode it to a pair (tx,ty)
3) read v1 from img[tx,ty]
4) do some calculations to extract a boolean value C and a target value T from v1,v,and the neighbours of (x,y) and (tx,ty)
5) *** IF ( C ) THEN WRITE T INTO (tx,ty).
Q1:所有“ C”为真的内核都将竞争写操作。假设哪一个获胜(最后写入)无关紧要。我已经在intel GPU上进行了一些测试,并且(禁用了过滤功能,并且启用了钳位功能)似乎没有任何问题,write_image似乎是原子性的,具有获胜的价值,并且我的算法收敛非常快。我可以安全地认为“未过滤”图像上的write_image是原子的吗?
Q2:我真正需要的是写入(tx,ty)从每个内核获得的最大T。那将涉及使用缓冲区而不是图像,夹紧自己(或使用更大的缓冲区填充零),并**在每个内核中使用atomic_max **。我还没有出于懒惰而这样做,因为我需要更改代码以仅使用缓冲区来测试它,但是我相信它会慢得多。我说的对吗?
为完整起见,这是我的实际内核(有待优化,欢迎提出任何建议!)
```
__kernel void color_components2(/* base image */ __read_only image2d_t image,/* uint32 */ __read_only image2d_t inputImage1,__write_only image2d_t outImage1) {
int2 gid = (int2)(get_global_id(0),get_global_id(1));
int x = gid.x;
int y = gid.y;
int lock = 0;
int2 size = get_image_dim(inputImage1);
const sampler_t sampler =
CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
uint4 base = read_imageui(image,sampler,gid);
uint4 ui4a = read_imageui(inputImage1,gid);
int2 t = (int2)(ui4a[0] % size.x,ui4a[0] / size.x);
unsigned int m = ui4a[0];
unsigned int n = ui4a[0];
if (base[0] > 0) {
for (int a = -1; a <= 1; a++)
for (int b = -1; b <= 1; b++) {
uint4 tmpa =
read_imageui(inputImage1,(int2)(t.x + a,t.y + b));
m = max(tmpa[0],m);
uint4 tmpb = read_imageui(inputImage1,(int2)(x + a,y + b));
n = max(tmpb[0],n);
}
}
if(n > m) write_imageui(outImage1,t,(uint4)(n,0));
}
```
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。