59

cudaMalloc()我对如何使用和分配和复制线性内存有很好的理解cudaMemcpy()。然而，当我想使用 CUDA 函数来分配和复制 2D 或 3D 矩阵时，我经常被各种参数弄糊涂，尤其是在处理 2D/3D 数组时总是存在的倾斜指针。该文档很好地提供了一些关于如何使用它们的示例，但它假设我熟悉填充和间距的概念，而我并不熟悉。

我通常最终会调整我在文档或网络上其他地方找到的各种示例，但随后的盲目调试非常痛苦，所以我的问题是：

什么是音高？我该如何使用它？如何在 CUDA 中分配和复制 2D 和 3D 数组？

4

1 回答 1

112

这是关于 CUDA 中的倾斜指针和填充的解释。

线性内存与填充内存

首先，让我们从非线性内存存在的原因说起。当使用 cudaMalloc 分配内存时，结果就像使用 malloc 分配一样，我们有一个指定大小的连续内存块，我们可以在其中放入任何我们想要的东西。如果我们想分配一个 10000 浮点数的向量，我们只需这样做：

float* myVector;
cudaMalloc(&myVector, 10000*sizeof(float));

然后通过经典索引访问 myVector 的第 i 个元素：

float element = myVector[i];

如果我们想访问下一个元素，我们只需：

float next_element = myvector[i+1];

它工作得很好，因为访问第一个元素旁边的元素（出于我不知道的原因，我现在不希望这样做）很便宜。

当我们将内存用作二维数组时，情况会有所不同。假设我们的 10000 浮点向量实际上是一个 100x100 数组。我们可以使用相同的 cudaMalloc 函数来分配它，如果我们想读取第 i 行，我们这样做：

float* myArray;
cudaMalloc(&myArray, 10000*sizeof(float));
int row[100];  // number of columns
for (int j=0; j<100; ++j)
    row[j] = myArray[i*100+j];

字对齐

所以我们必须从 myArray+100*i 读取内存到 myArray+101*i-1。它将执行的内存访问操作数取决于该行占用的内存字数。内存字中的字节数取决于实现。为了在读取单行时最小化内存访问次数，我们必须确保我们从单词的开头开始行，因此我们必须为每一行填充内存，直到新行的开头。

银行冲突

填充数组的另一个原因是 CUDA 中的存储库机制，涉及共享内存访问。当数组在共享内存中时，它被分成几个内存库。两个 CUDA 线程可以同时访问它，前提是它们不访问属于同一内存库的内存。由于我们通常希望并行处理每一行，因此我们可以通过将每一行填充到新银行的开头来确保可以模拟访问它。

现在，我们将使用 cudaMallocPitch，而不是使用 cudaMalloc 分配 2D 数组：

size_t pitch;
float* myArray;
cudaMallocPitch(&myArray, &pitch, 100*sizeof(float), 100);  // width in bytes by height

请注意，这里的音高是函数的返回值：cudaMallocPitch 检查它在您的系统上应该是什么并返回适当的值。cudaMallocPitch 的作用如下：

  1. 分配第一行。
  2. 检查分配的字节数是否使其正确对齐。例如，它是 128 的倍数。
  3. 如果不是，则分配更多字节以达到 128的下一个倍数。间距是为单行分配的字节数，包括额外字节（填充字节）。
  4. 重申每一行。

最后，我们通常分配了比需要更多的内存，因为现在每一行都是 pitch 的大小，而不是w*sizeof(float).

但是现在，当我们想要访问列中的元素时，我们必须这样做：

float* row_start = (float*)((char*)myArray + row * pitch);
float column_element = row_start[column];

两个连续列之间的字节偏移量不能再从我们数组的大小中推断出来，这就是为什么我们要保持 cudaMallocPitch 返回的音高。而且由于音高是填充大小的倍数（通常是字大小和库大小中的最大值），所以效果很好。耶。

将数据复制到/从音高内存

现在我们知道如何创建和访问由 cudaMallocPitch 创建的数组中的单个元素，我们可能希望将它的整个部分复制到其他内存或从其他内存中复制，无论是否是线性的。

假设我们想用 malloc 将我们的数组复制到我们的主机上分配的 100x100 数组中：

float* host_memory = (float*)malloc(100*100*sizeof(float));

如果我们使用 cudaMemcpy，我们将复制使用 cudaMallocPitch 分配的所有内存，包括每行之间的填充字节。为了避免填充内存，我们必须做的是逐行复制。我们可以手动完成：

for (size_t i=0; i<100; ++i) {
  cudaMemcpy(host_memory[i*100], myArray[pitch*i],
             100*sizeof(float), cudaMemcpyDeviceToHost);
}

或者我们可以告诉 CUDA API，为了方便起见，我们只需要使用填充字节分配的内存中的有用内存，所以如果它可以自动处理自己的混乱，那确实非常好，谢谢。这里进入 cudaMemcpy2D：

cudaMemcpy2D(host_memory, 100*sizeof(float)/*no pitch on host*/,
             myArray, pitch/*CUDA pitch*/,
             100*sizeof(float)/*width in bytes*/, 100/*heigth*/, 
             cudaMemcpyDeviceToHost);

现在复制将自动完成。它将复制宽度（此处：100xsizeof(float)）、高度时间（此处：100）、跳过间距指定的字节数每次跳转到下一行时的字节数。请注意，我们仍然必须为目标内存提供音高，因为它也可以被填充。这里不是，所以间距等于非填充数组的间距：它是一行的大小。另请注意，memcpy 函数中的宽度参数以字节表示，但高度参数以元素数量表示。那是因为复制的方式，有点像我在上面写的手动复制：宽度是每个副本沿一行的大小（在内存中连续的元素），高度是这个操作必须的次数得以实现。（作为物理学家，这些单位的不一致让我非常恼火。）

处理 3D 数组

3D 数组实际上与 2D 数组没有什么不同，没有包含额外的填充。3D 数组只是填充行的 2D经典数组。这就是为什么在分配 3D 数组时，您只会得到一个音高，即沿行的连续点之间的字节数差异。如果您想访问沿深度维度的连续点，您可以安全地将间距乘以列数，从而得到 slicePitch。

用于访问 3D 内存的 CUDA API 与用于访问 2D 内存的 API 略有不同，但思路是相同的：

  • 使用 cudaMalloc3D 时，您会收到一个音高值，您必须小心保留该值，以便后续访问内存。
  • 复制 3D 内存块时，除非复制单行，否则不能使用 cudaMemcpy。您必须使用考虑音高的 CUDA 实用程序提供的任何其他类型的复制实用程序。
  • 当您将数据复制到/从线性内存中时，您必须为指针提供一个间距，即使它无关紧要：这个间距是行的大小，以字节表示。
  • 大小参数以字节表示行大小，以元素数表示列和深度维度。
于 2013-04-20T11:43:33.153 回答