CUDA - memcpy2d - 错误的音高

Question

我刚刚开始 CUDA 编程，并尝试执行下面所示的代码。这个想法是将一个二维数组复制到设备，计算所有元素的总和，然后检索总和（我知道这个算法不是并行的。事实上，它正在做更多的工作，然后是必要的。然而，这只是有意的作为内存复制的练习）。

#include<stdio.h>
#include<cuda.h>
#include <iostream>
#include <cutil_inline.h>

#define height 50
#define width 50

using namespace std;

// Device code
__global__ void kernel(float* devPtr, int pitch,int* sum)
{
int tempsum = 0;    
for (int r = 0; r < height; ++r) {
        int* row = (int*)((char*)devPtr + r * pitch);
        for (int c = 0; c < width; ++c) {
             int element = row[c];
             tempsum = tempsum + element;
        }
    }
*sum = tempsum;
}

//Host Code
int main()
{

int testarray[2][8] = {{4,4,4,4,4,4,4,4},{4,4,4,4,4,4,4,4}};
int* sum =0;
int* sumhost = 0;
sumhost = (int*)malloc(sizeof(int));

cout << *sumhost << endl;

float* devPtr;
size_t pitch;
cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(int), height);
cudaMemcpy2D(devPtr,pitch,testarray,0,8* sizeof(int),4,cudaMemcpyHostToDevice);

cudaMalloc((void**)&sum, sizeof(int));
kernel<<<1, 4>>>(devPtr, pitch, sum);
cutilCheckMsg("kernel launch failure");
cudaMemcpy(sumhost, sum, sizeof(int), cudaMemcpyDeviceToHost);

cout << *sumhost << endl;

return 0;
}

这段代码编译得很好（在 4.0 sdk 候选版本上）。然而，一旦我尝试执行，我就会得到

0
cpexample.cu(43) : cutilCheckMsg() CUTIL CUDA error : kernel launch failure : invalid pitch argument.

不幸的是，因为我不知道如何修复它；-(。据我所知，音调是内存中的偏移量，以允许更快地复制数据。但是这样的音调仅在设备内存中使用，而不是在设备内存中使用。在主机内存中，不是吗？因此我的主机内存的音高应该是0，不是吗？

另外我还想问另外两个问题：

• 如果我声明一个像 int* sumhost 这样的变量（见上文），那么这个指针指向哪里？首先到主机内存，然后cudaMalloc到设备内存？
• cutilCheckMsg 在这种情况下非常方便。我应该知道类似的调试功能吗？

Answer 1

在您的这行代码中：

cudaMemcpy2D(devPtr,pitch,testarray,0,8* sizeof(int),4,cudaMemcpyHostToDevice);

您是说 的源音高值testarray等于0，但是当音高的公式为 时，这怎么可能呢T* elem = (T*)((char*)base_address + row * pitch) + column？如果我们在该公式中替换pitch 的值0，则在某个二维(x,y) 有序对偏移处查找地址时，我们将无法获得正确的值。需要考虑的一件事是音高值的规则是pitch = width + padding。在主机上，填充通常等于0，但宽度则不然0，除非数组中没有任何内容。在硬件方面可能有额外的填充，这就是为什么间距值可能不等于数组的声明宽度的原因。pitch >= width因此，您可以根据填充值得出结论。因此，即使在主机端，源间距的值也应至少为每行的大小（以字节为单位），这意味着在 的情况下testarray，它应该是8*sizeof(int)。最后，主机中二维数组的高度也只是2rows，而不是4。

作为关于分配指针会发生什么的问题的答案，如果您使用 分配指针malloc()，则该指针将被赋予驻留在主机内存中的地址值。因此，您可以在主机端取消引用它，但不能在设备端取消引用。另一方面，用 分配的指针cudaMalloc()被赋予指向驻留在设备上的内存的指针。因此，如果您在主机上取消引用它，它不会指向主机上分配的内存，并且会发生不可预测的结果。不过，可以将此指针地址传递给设备上的内核，因为当它在设备端取消引用时，它指向设备本地可访问的内存。总体而言，CUDA 运行时将这两个内存位置分开，提供内存复制功能，在设备和主机之间来回复制，并根据所需方向使用这些指针的地址值作为复制的源和/或目标（主机到设备或设备到主机）。现在，如果您采用相同的int*，并首先使用 分配它malloc()，然后（希望在调用free()指针之后）使用cudaMalloc()，您的指针将首先有一个指向主机内存的地址，然后是设备内存。您必须跟踪其状态，以避免因取消引用设备或主机上的地址而产生不可预测的结果，具体取决于它是在主机代码还是设备代码中取消引用。