c++ - cudaMemcpy 段错误

Question

我被这个错误困扰了很长一段时间，所以我决定把它贴在这里。

调用 cudaMemcpy 时发生此段错误:

CurrentGrid->cdata[i] = new float[size];
cudaMemcpy(CurrentGrid->cdata[i], Grid_dev->cdata[i], size*sizeof(float),\
                cudaMemcpyDeviceToHost);

CurrentGrid 和 Grid_dev 分别是指向主机和设备上的 grid 类对象的指针，在此上下文中 i=0。类成员cdata是一个float类型的指针数组。为了调试，就在这个 cudaMemcpy 调用之前，我打印出了 Grid_Dev->cdata[i] 的每个元素的值，CurrentGrid->cdata[i] 的地址和Grid_dev->cdata[i] 和 size 的值，看起来都不错。但它仍然以“Segmentation fault (core dumped)”结束，这是唯一的错误消息。 cuda-memcheck 只给出了“进程没有成功终止”。我现在无法使用 cuda-gdb。关于去哪里有什么建议吗？

更新:现在看来我已经通过 cudaMalloc 解决了这个问题，方法是设备上的另一个浮点指针 A 和 cudaMemcpy Grid_dev->cdata[i] 的值到 A，然后 cudaMemcpy A 到主机。所以上面写的这段代码就变成了:

float * A;
cudaMalloc((void**)&A, sizeof(float));
...
...
cudaMemcpy(&A, &(Grid_dev->cdata[i]), sizeof(float *), cudaMemcpyDeviceToHost);    
CurrentGrid->cdata[i] = new float[size];
cudaMemcpy(CurrentGrid->cdata[i], A, size*sizeof(float), cudaMemcpyDeviceToHost);            

我这样做是因为 valgrind 弹出“大小为 8 的无效读取”，我认为它指的是 Grid_dev->cdata[i]。我用 gdb 再次检查，打印出 Grid_dev->cdata[i] 的值为 NULL。所以我想即使在这个 cudaMemcpy 调用中我也不能直接取消引用设备指针。但为什么 ？根据这个thread底部的评论，我们应该能够在 cudaMemcpy 函数中取消引用设备指针。

此外，我不知道 cudaMalloc 和 cudaMemcpy 如何工作的底层机制，但我认为通过 cudaMalloc 一个指针，这里说 A，我们实际上将这个指针分配给指向设备上的某个地址。并且通过cudaMemcpy将Grid_dev->cdata[i]修改为上面代码中的A，我们重新赋值指针A指向数组。那我们在cudaMalloced的时候不就失去了A指向的前一个地址的踪迹了吗？这会导致内存泄漏或其他什么吗？如果是，我应该如何正确解决这种情况？谢谢!

为了引用，我把发生这个错误的完整函数的代码放在下面。

非常感谢!

__global__ void Print(grid *, int);
__global__ void Printcell(grid *, int);
void CopyDataToHost(param_t p, grid * CurrentGrid, grid * Grid_dev){

    cudaMemcpy(CurrentGrid, Grid_dev, sizeof(grid), cudaMemcpyDeviceToHost);
#if DEBUG_DEV
    cudaCheckErrors("cudaMemcpy1 error");
#endif
    printf("\nBefore copy cell data\n");
    Print<<<1,1>>>(Grid_dev, 0);            //Print out some Grid_dev information for 
    cudaDeviceSynchronize();                //debug 
    int NumberOfBaryonFields = CurrentGrid->ReturnNumberOfBaryonFields();
    int size = CurrentGrid->ReturnSize();
    int vsize = CurrentGrid->ReturnVSize();
    CurrentGrid->FieldType = NULL;
    CurrentGrid->FieldType = new int[NumberOfBaryonFields];
    printf("CurrentGrid size is %d\n", size);
    for( int i = 0; i < p.NumberOfFields; i++){
        CurrentGrid->cdata[i] = NULL;
        CurrentGrid->vdata[i] = NULL;
        CurrentGrid->cdata[i] = new float[size];
        CurrentGrid->vdata[i] = new float[vsize];

        Printcell<<<1,1>>>(Grid_dev, i);//Print out element value of Grid_dev->cdata[i]
        cudaDeviceSynchronize();        

        cudaMemcpy(CurrentGrid->cdata[i], Grid_dev->cdata[i], size*sizeof(float),\
                cudaMemcpyDeviceToHost);               //where error occurs
#if DEBUG_DEV
        cudaCheckErrors("cudaMemcpy2 error");
#endif
        printf("\nAfter copy cell data\n");
        Print<<<1,1>>>(Grid_dev, i);
        cudaDeviceSynchronize();
        cudaMemcpy(CurrentGrid->vdata[i], Grid_dev->vdata[i], vsize*sizeof(float),\
                cudaMemcpyDeviceToHost);
#if DEBUG_DEV
        cudaCheckErrors("cudaMemcpy3 error");
#endif
    }
    cudaMemcpy(CurrentGrid->FieldType, Grid_dev->FieldType,\
            NumberOfBaryonFields*sizeof(int), cudaMemcpyDeviceToHost);
#if DEBUG_DEV
    cudaCheckErrors("cudaMemcpy4 error");
#endif
}

编辑:这是来自 valgrind 的信息，我正试图从中找出发生内存泄漏的位置。

==19340== Warning: set address range perms: large range [0x800000000, 0xd00000000) (noaccess)
==19340== Warning: set address range perms: large range [0x200000000, 0x400000000) (noaccess)
==19340== Invalid read of size 8
==19340==    at 0x402C79: CopyDataToHost(param_t, grid*, grid*) (CheckDevice.cu:48)
==19340==    by 0x403646: CheckDevice(param_t, grid*, grid*) (CheckDevice.cu:186)
==19340==    by 0x40A6CD: main (Transport.cu:81)
==19340==  Address 0x2003000c0 is not stack'd, malloc'd or (recently) free'd
==19340== 
==19340== 
==19340== Process terminating with default action of signal 11 (SIGSEGV)
==19340==  Bad permissions for mapped region at address 0x2003000C0
==19340==    at 0x402C79: CopyDataToHost(param_t, grid*, grid*) (CheckDevice.cu:48)
==19340==    by 0x403646: CheckDevice(param_t, grid*, grid*) (CheckDevice.cu:186)
==19340==    by 0x40A6CD: main (Transport.cu:81)
==19340== 
==19340== HEAP SUMMARY:
==19340==     in use at exit: 2,611,365 bytes in 5,017 blocks
==19340==   total heap usage: 5,879 allocs, 862 frees, 4,332,278 bytes allocated
==19340== 
==19340== LEAK SUMMARY:
==19340==    definitely lost: 0 bytes in 0 blocks
==19340==    indirectly lost: 0 bytes in 0 blocks
==19340==      possibly lost: 37,416 bytes in 274 blocks
==19340==    still reachable: 2,573,949 bytes in 4,743 blocks
==19340==         suppressed: 0 bytes in 0 blocks
==19340== Rerun with --leak-check=full to see details of leaked memory
==19340== 
==19340== For counts of detected and suppressed errors, rerun with: -v
==19340== ERROR SUMMARY: 1 errors from 1 contexts (suppressed: 2 from 2)

Answer 1

我相信我知道问题出在哪里，但为了确认它，查看您用于在设备上设置 Grid_dev 类的代码会很有用。

当一个类或其他数据结构要在设备上使用时，该类中有指针指向内存中的其他对象或缓冲区(可能在设备内存中，用于将在设备上使用的类)，那么让这个顶级类在设备上可用的过程就变得更加复杂。

假设我有这样一个类:

class myclass{
  int myval;
  int *myptr;
  }

我可以在主机上实例化上面的类，然后 malloc 一个 int 的数组并将该指针分配给 myptr，等等会没事的。要使此类在设备上可用且仅在设备上可用，过程可能类似。我可以:

1. cudaMalloc 指向将保存 myclass 的设备内存的指针>
2. (可选)使用 cudaMemcpy 将主机上 myclass 的实例化对象复制到步骤 1 中的设备指针
3. 在设备上，使用malloc或new为myptr分配设备存储空间

如果我不想访问主机上为 myptr 分配的存储空间，则上述顺序很好。但是，如果我确实希望该存储从主机可见，我需要一个不同的顺序:

1. cudaMalloc 指向将保存 myclass 的设备内存的指针，我们称之为 mydevobj
2. (可选)使用 cudaMemcpy 将主机上 myclass 的实例化对象复制到设备指针 mydevobj
3. 在主机上创建一个单独的 int 指针，我们称它为 myhostptr
4. cudaMalloc int 存储在设备上，用于 myhostptr
5. cudaMemcpy 主机myhostptr的指针值到设备指针&(mydevobj->myptr)

之后，你可以cudaMemcpy 嵌入指针myptr指向的数据到上分配的区域(通过cudaMalloc) >我的主机名

注意，在第5步中，因为我取的是这个指针位置的地址，所以这个cudaMemcpy操作只需要主机上的mydevobj指针，这在一个cudaMemcpy操作中有效(仅)。

然后设备指针 myint 的值将被正确设置以执行您尝试执行的操作。如果您随后想要将 cudaMemcpy 数据传入和传出 myint 到主机，您可以在任何 cudaMemcpy 调用中使用指针 myhostptr，不是 mydevobj->myptr.如果我们尝试使用 mydevobj->myptr，则需要取消引用 mydevobj，然后使用它来检索存储在 myptr 中的指针，然后使用该指针作为到/从位置的拷贝。这在主机代码中是 Not Acceptable 。如果您尝试这样做，您将遇到段错误。 (请注意，通过类比，我的 mydevobj 就像您的 Grid_dev 而我的 myptr 就像您的 cdata)

总的来说，这是一个第一次遇到它时需要仔细考虑的概念，因此像这样的问题在 SO 上出现了一些频率。您可能想研究其中一些问题以查看代码示例(因为您尚未提供设置 Grid_dev 的代码):

1. example 1
2. example 2
3. example 3