我应该如何以及何时使用cuda API使用倾斜指针？

Question

我对如何使用cudaMalloc()和分配和复制线性内存有很好的理解cudaMemcpy().但是,当我想使用CUDA函数来分配和复制2D或3D矩阵时,我常常会被各种参数所迷惑,特别是关于在处理2D/3D数组时总是存在的倾斜指针.文档很适合提供一些如何使用它们的例子,但它假设我熟悉填充和音高的概念,我不是.

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

什么是球场？我该如何使用它？如何在CUDA中分配和复制2D和3D阵列？

Answer 1

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

线性内存与填充内存

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

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

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

float element = myVector[i];

如果我们想要访问下一个元素,我们只需:

float next_element = myvector[i+1];

它工作得非常好,因为访问第一个元素旁边的元素(因为我不知道并且我不希望现在的原因)便宜.

当我们将内存用作2D数组时,情况会有所不同.假设我们的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. 重申每一行.

最后,我们通常会分配比所需更多的内存,因为每行现在都是音高的大小,而不是大小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函数中的width参数以字节表示,但height参数以元素数表示.这是因为副本的完成方式,就像我上面写的手册一样:宽度是一行中每个副本的大小(内存中连续的元素),高度是此操作必须的次数完成.(作为一名物理学家,这些单位的不一致让我非常恼火.)

处理3D阵列

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

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

• 使用cudaMalloc3D时,您会收到一个音高值,您必须小心保留该值以便随后访问内存.
• 复制3D内存块时,除非复制单行,否则不能使用cudaMemcpy.您必须使用CUDA实用程序提供的任何其他类型的复制实用程序.
• 将数据复制到线性存储器或从线性存储器复制数据时,即使它与指针无关,也必须为指针提供间距:此间距是行的大小,以字节为单位表示.
• 大小参数以行大小的字节数表示,并以列和深度维度的元素数量表示.