c - Cuda 用 2D 块替换双倍

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

撇开数组分配和索引方案是否正确（我不确定帖子中是否有足够的信息来证实这一点），以及整数除法和模数很慢并且应该避免的事实，你有一个更根本的问题 - 内存竞赛。

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