除了方便之外，在 CUDA 中使用 2d 内核还有其他优势吗？

Question

我想知道除了抽象的便利之外，让内核 > 1d 是否有任何内在优势。我认为如果内核的维度相关，那么答案可能与 gpu 的布局有关。我通常更愿意坚持 1d 并展平更高维度的数据。这种做法在技术层面有什么问题吗？

Answer 1

两者在性能方面可能不等同：展平 2D/3D 位置很便宜，但是从 1D 展平位置计算 2D 或 3D 块位置是昂贵的，因为这需要一个缓慢的 modulus/division（块并不总是 power of 也不是编译时已知的）。更不用说使用 2D 网格进行 2D 计算可以使代码更具可读性（意图更明确）。对于 3D 内核尤其如此。而且，dimensions are bounded。如果您想对大于 2 GB 的二维数组执行计算，这可能是一个限制。

Answer 2

更糟糕的是 Jerome 所说的...结合 2D/3D 位置没有那么便宜。想想看：

flattened_block_id = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.y * gridDim.z;
flattened_thread_id = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.y * blockDim.z; 
block_volume = blockDim.x * blockDim.y * blockDim.z;
global_flattened_id = flattened_thread_id + block_volume * flattened_block_id;

这就是忽略维度边界时的全部内容。一般来说，你不能这样做。因此，您需要一些符号扩展，并且一半的乘法和加法变为 64 位。这是很多操作！想想那些你有条件的情况，比如：

if (is_nice(global_flattened_id)) { return; }

有了它，您只需确保您必须为所有这些操作付费，即使您的线程不会执行任何操作。

话虽如此...当我编写处理一维数据的内核时，我也认为这些额外的维度很愚蠢。然后我开始真正拥有 3D（或 5D）数据，它们很快就派上了用场:-)

最后，请记住：CUDA 是 3D graphics shader evolution 的副产品。在需要您或我作为用户之前，它需要 3D 表示...

除了方便之外，在 CUDA 中使用 2d 内核还有其他优势吗？

Is there an advantage to using 2d kernels in CUDA beyond convenience?

c++

cuda

dimensions