【发布时间】：2021-10-29 10:01:23
【问题描述】：
我想知道除了抽象的便利性之外，让内核大于 1d 是否有任何固有的优势。我认为如果内核的尺寸相关，答案可能与 gpu 的布局有关。我通常更喜欢坚持一维并展平更高维度的数据。这种方法在技术层面有什么问题吗？
【问题讨论】：
标签： c++ cuda dimensions
【发布时间】：2021-10-29 10:01:23
【问题描述】：
两者在性能方面可能不相等：展平 2D/3D 位置很便宜，但从 1D 展平位置计算 2D 或 3D 块位置很昂贵，因为这需要缓慢的模数/除法（块并不总是的力量在编译时也不知道）。更不用说使用 2D 网格进行 2D 计算使代码更具可读性（意图更清晰）。对于 3D 内核尤其如此。此外，dimensions are bounded。如果您想在大于 2 GB 的二维数组上执行计算，这可能是一个限制。
【讨论】：
更糟糕的是 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 表示...
【讨论】：
-
感谢您提供信息丰富的回复。后续问题：尺寸的大小可以是任意的，还是我应该特别注意 gpu 的实际布局（即将内核初始化为我的问题的尺寸，或者使我的问题符合内核的最佳尺寸)?我知道这是一个广泛的问题。您是否知道任何有助于阐明 gpu 架构和内核维度之间关系的资源（尤其是视频）？再次感谢
-
@classic_sasquatch_behavior：尺寸有不同的尺寸限制。 x 是“大”维度，y 和 z 是“小”；见here。
-
虽然像
blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.y * gridDim.z这样的行在视觉上看起来很糟糕，但它们的计算量并不那么大，因为 GPU 通常可以使用像MAD这样的 基于 FMA 的指令 来非常快速地计算它. Here 就是一个例子。请注意，这两种实现方式的移动次数都非常高。第一个实现只使用了一些额外的相对便宜的指令。除此之外，我同意边界和其余部分；）。 -
@JérômeRichard：您的示例仅适用于 32 位的情况。对于 64 位，它看起来像 this。
mov用于获取特殊的寄存器值（这本身并不像您期望的那么便宜）。 -
即使对于巨大的数组，我希望
flattened_block_id不会在 32 位中溢出，因为每个块至少应该有 32 个线程（为了性能）。如果发生溢出，则意味着数组大于 64 GiB，这太大而无法放入 GPU 内存（并且使用统一内存会贵得多）。只有global_flattened_id需要是 64 位的（就像在其他内核中一样）。对于mov，您是否知道低级程序集是从专用寄存器、RAM（缓存）变量还是从任何其他数据源获取它们的？如果他们也可以被驱逐，IDK。