CUDA,使用 2D 和 3D 数组

Posted

技术标签:

【中文标题】CUDA,使用 2D 和 3D 数组【英文标题】:CUDA, Using 2D and 3D Arrays 【发布时间】:2018-01-20 11:19:46 【问题描述】:

网上有很多关于在 CUDA 上分配、复制、索引等 2d 和 3d 数组的问题。我得到了很多相互矛盾的答案,所以我试图编译过去的问题,看看我是否可以问正确的问题。

第一个链接:https://devtalk.nvidia.com/default/topic/392370/how-to-cudamalloc-two-dimensional-array-/

问题:分配一个二维指针数组

用户解决方案:使用 mallocPitch

“正确”的低效解决方案:在 for 循环中为每一行使用 malloc 和 memcpy(荒谬的开销)

“更正确”的解决方案:将其压缩成一维数组“专业意见”,一条评论说没有人关注性能在 gpu 上使用二维指针结构

第二个链接:https://devtalk.nvidia.com/default/topic/413905/passing-a-multidimensional-array-to-kernel-how-to-allocate-space-in-host-and-pass-to-device-/

问题:在主机上分配空间并将其传递给设备

子链接:https://devtalk.nvidia.com/default/topic/398305/cuda-programming-and-performance/dynamically-allocate-array-of-structs/

子链接解决方案:在 GPU 上编码基于指针的结构是一种糟糕的体验并且效率非常低,将其压缩成一维数组。

第三个链接:Allocate 2D Array on Device Memory in CUDA

问题:分配和传输二维数组

用户解决方案:使用 mallocPitch

其他解决方案:将其展平

第四个链接:How to use 2D Arrays in CUDA?

问题:分配和遍历二维数组

提交的解决方案:不显示分配

其他解决方案:压扁它

有很多其他消息来源大多都在说同样的事情,但在多个实例中,我看到有关 GPU 上的指针结构的警告。

许多人声称分配指针数组的正确方法是为每一行调用 malloc 和 memcpy,但存在 mallocPitch 和 memcpy2D 函数。这些功能是否效率较低?为什么这不是默认答案?

二维数组的另一个“正确”答案是将它们压缩成一个数组。我应该习惯这作为生活的事实吗?我对我的代码非常挑剔,我觉得它不雅。

我正在考虑的另一个解决方案是最大化使用一维指针数组的矩阵类,但我找不到实现双括号运算符的方法。

也可以根据这个链接:Copy an object to device?

和子链接回答:cudaMemcpy segmentation fault

这有点不确定。

我想使用 CUDA 的类都有 2/3d 数组,将这些类转换为 CUDA 的 1d 数组不会有很多开销吗?

我知道我已经问了很多,但总而言之,我是否应该习惯将压缩数组作为生活中的事实,或者我是否可以使用 2d 分配和复制函数而不会像在调用 alloc 和 cpy 的解决方案中那样产生糟糕的开销在for循环中?

【问题讨论】:

【参考方案1】:

由于您的问题汇总了其他问题的列表,因此我将通过汇总其他答案的列表来回答。

cudaMallocPitch/cudaMemcpy2D:

首先,像 cudaMallocPitchcudaMemcpy2D 这样的 cuda 运行时 API 函数实际上不涉及双指针分配或 2D(双下标)数组。这很容易通过查看the documentation 并注意函数原型中的参数类型来确认。 srcdst 参数是单指针参数。它们不能被双重下标或双重取消引用。对于其他示例用法,here 是有关此问题的众多问题之一。 here 是一个完整的示例用法。另一个涵盖与cudaMallocPitch/cudaMemcpy2d 用法相关的各种概念的示例是here。相反,考虑这些的正确方法是它们使用 pitched 分配。此外,在循环中使用一组malloc(或new,或类似的)操作创建基础分配时,您不能使用cudaMemcpy2D 传输数据。这种主机数据分配结构特别不适合处理设备上的数据。

一般的、动态分配的 2D 案例:

如果您想了解如何在 CUDA 内核中使用动态分配的二维数组(这意味着您可以使用双下标访问,例如data[x][y]),那么cuda tag info page 包含对此的“规范”问题,它是here。 talonmies 给出的答案包括适当的机制以及适当的警告:

还有额外的、非平凡的复杂性 访问的效率通常低于一维访问,因为数据访问需要取消引用 2 个指针,而不是 1 个。

(请注意,分配对象数组,其中对象具有指向动态分配的嵌入式指针,本质上与二维数组概念相同,您在问题中链接的the example 是一个合理的演示为此)

另外,here 是一种用于构建通用动态分配二维数组的推力方法。

扁平化:

如果你认为你必须使用一般的 2D 方法,那就继续吧,这不是不可能的(虽然有时people struggle 有这个过程!)但是,由于增加了复杂性和降低了效率,这里的规范“建议”是“扁平化”您的存储方法,并使用“模拟”2D 访问。 Here 是许多讨论“扁平化”的问题/答案示例之一。

一般、动态分配的 3D 案例:

当我们将其扩展到 3(或更高!)维度时,IMO 处理一般情况变得过于复杂。额外的复杂性应该强烈地激励我们寻求替代方案。三下标的一般情况在实际检索数据之前涉及 3 次指针访问,因此效率更低。 Here 是一个完整的示例(第二个代码示例)。

特殊情况:编译时已知的数组宽度:

请注意,当数组维度(宽度,在二维数组的情况下,或 3 个中的 2 个)时,应将其视为特殊情况 3D 数组的尺寸)在编译时是已知的。在这种情况下,通过适当的辅助类型定义,我们可以“指示”编译器应该如何计算索引,并且在这种情况下,我们可以使用双下标访问,其复杂性远低于一般情况 和 不会因为指针追逐而损失效率。只需要取消引用一个指针即可检索数据(无论数组维数如何,如果在编译时 n 维数组的 n-1 维是已知的)。已经提到的答案here(第一个代码示例)中的第一个代码示例给出了 3D 情况下的完整示例,答案here 给出了这种特殊情况的 2D 示例。

双下标主机代码、单下标设备代码:

最后,另一个方法选项允许我们在主机代码中轻松混合 2D(双下标)访问,同时在 中仅使用 1D(单下标,可能带有“模拟 2D”访问)设备代码。一个可行的例子是here。通过将底层分配组织为连续分配,然后构建指针“树”,我们可以在主机上启用双下标访问,并且仍然可以轻松地将平面分配传递给设备。尽管示例未显示,但可以扩展此方法以基于平面分配和手动创建的指针“树”在设备上创建双下标访问系统,但这将具有大致相同的问题正如上面给出的 2D 通用动态分配方法:它将涉及双指针(双解引用)访问,因此效率较低,并且与构建指针“树”相关,以用于设备代码(例如,它会可能需要额外的cudaMemcpy 操作)。

从上述方法中,您需要选择一种适合您的胃口和需求的方法。没有一个建议适合所有可能的情况。

【讨论】:

以上是关于CUDA,使用 2D 和 3D 数组的主要内容,如果未能解决你的问题,请参考以下文章

将 3d 数组发送到 CUDA 内核

使用Cuda平行降维(3D到2D,总和)

我应该如何以及何时将倾斜指针与 cuda API 一起使用?

二维数组 cudaMallocPitch() 和三维数组 cudaMalloc3D() 的使用

如何使用CUDA并行化嵌套for循环以在2D数组上执行计算

NumPy:在 3D 切片中使用来自 argmin 的 2D 索引数组