Cuda 用 2D block 替换 double

Question

我对 CUDA 很陌生，一直在尝试遍历 2D 数组。我有以下代码，可以在纯 C 上按预期工作:

for (ty=0;ty<s;ty++){
        if (ty+pixY < s && ty+pixY>=0){
            for(tx=0;tx<r;tx++){
                T[ty/3][tx/3] += (tx+pixX<s && tx+pixX>=0) ? 
                *(image +M*(ty+pixY)+tx+pixX) * *(filter+fw*(ty%3)+tx%3) : 0;
            }
        }
    }  

也许我弄错了，但是这段代码不会像下面这样转换为 CUDA 吗？

tx = threadIdx.x;
ty = threadIdy.y;

T[ty/3][tx/3] += (tx+pixX<s && tx+pixX>=0) ?
                *(image +M*(ty+pixY)+tx+pixX) * *(filter+fw*(ty%3)+tx%3) : 0; 

假设我已将内核参数定义为 dimGrid(1,1,1) 和 blockDim(r,s,1)

我问是因为我得到了意想不到的结果。另外，如果我正确地将数组声明并分配为 2D cuda 数组，而不仅仅是一个大的 1D 数组，这会有帮助吗？

感谢您的帮助。

Answer 1

抛开数组分配和索引方案是否正确(我不确定帖子中是否有足够的信息来确认这一点)，以及整数除法和取模很慢并且应该避免的事实，您有很多更根本的问题 - 内存竞争。

您正在使用的单个 block 中的多个线程将尝试同时读取和写入 T 的同一条目。 CUDA 不保证此类操作的正确性，而且几乎可以肯定它不会起作用。最简单的替代方案是仅使用单个线程来计算每个 T[][] 条目，而不是三个线程。这消除了内存竞争。