我正在编写一个简单的光线追踪器。这个想法是，对于每个像素，都有一个线程遍历驻留在全局内存中的某个结构（几何）。
我像这样调用我的内核：
trace<<<gridDim, blockDim>>>(width, height, frameBuffer, scene)
其中scene 是先前分配有cudaMalloc 的结构。每个线程都必须从同一个节点开始遍历这个结构，并且很可能许多并发线程会尝试多次读取同一个节点。这是否意味着当这样的读取发生时，它会削弱并行度？
鉴于几何体很大，我认为复制它不是一种选择。我的意思是整个处理过程仍然发生得相当快，但我想知道这是必须处理的事情，还是只是随随便便。
【问题讨论】：
标签： cuda gpgpu raytracing
首先，当您说并发读取可能会或可能不会削弱并行度时，我认为您的想法是错误的。因为这就是并行的含义。每个线程都在同时读取。相反，当每个线程基本上想要相同的东西，即相同的节点时，您应该考虑它是否会由于更多的内存访问而影响性能。
根据文章here，如果数据局部性存在且仅在warp 内，则可以合并内存访问。
这意味着如果 warp 中的线程试图访问彼此靠近的内存位置，它们可以被合并。在您的情况下，每个线程都试图访问“相同”节点，直到它遇到它们分支的端点。
这意味着内存访问将在 warp 中合并，直到线程分支。
【讨论】：
从每个线程高效访问全局内存取决于您的设备架构和代码。 CUDA 驱动程序将分配在全局内存上的数组与 256 字节内存段对齐。该设备可以通过与其大小对齐的 32、64 或 128 字节事务访问全局内存。该设备将warp线程发出的全局内存负载和存储合并到尽可能少的事务中，以最小化DRAM带宽。计算能力低于 2.0 的设备的数据访问未对齐会影响访问数据的有效带宽。当使用计算能力大于 2.0 的设备时，这不是一个严重的问题。话虽这么说，几乎不管您的设备是哪一代，当以较大的步幅访问全局内存时，有效带宽都会变差（Reference）。我会假设对于随机访问，可能会出现相同的行为。
【讨论】：
除非您在阅读时不更改结构，我假设您会这样做（如果这是一个场景，您可能会渲染每一帧？）然后是的，它会削弱性能并可能导致未定义的行为。这称为竞争条件。您可以使用原子操作来克服此类问题。使用原子操作可以保证不会发生竞争条件。 如果可以的话，您可以尝试将“场景”填充到共享内存中。 您还可以尝试使用流来增加并发性，这也会为在同一流中运行的内核带来某种同步。
【讨论】：