指向 CUDA 中对象的指针数组的指针

Question

我遵循了这个问题和这个链接提供的指导，这些指导涉及将指针数组传递到设备并返回的概念，但当指针指向一个对象时，我似乎正在努力解决我的特定情况。请参阅下面的示例代码，为简洁起见，我删除了错误检查。

// Kernel
__global__ void myKernel(Obj** d_array_of_objs)
{
    // Change the scalar of each object to 5
    // by dereferencing device array to get 
    // appropriate object pointer.
    *d_array_of_objs->changeToFive();    <--------- SEE QUESTION 4
}

// Entry point
int main()
{

    /********************************/
    /* INITIALISE OBJ ARRAY ON HOST */
    /********************************/

    // Array of 3 pointers to Objs
    Obj* h_obj[3];
    for (int i = 0; i < 3; i++) {
        h_obj[i] = new Obj();       // Create
        h_obj[i]->scalar = i * 10;  // Initialise
    }

    // Write out
    for (int i = 0; i < 3; i++) {
        std::cout << h_obj[i]->scalar << std::endl;
    }


    /**************************************************/
    /* CREATE DEVICE VERSIONS AND STORE IN HOST ARRAY */
    /**************************************************/

    // Create host pointer to array-like storage of device pointers
    Obj** h_d_obj = (Obj**)malloc(sizeof(Obj*) * 3);    <--------- SEE QUESTION 1
    for (int i = 0; i < 3; i++) {
        // Allocate space for an Obj and assign
        cudaMalloc((void**)&h_d_obj[i], sizeof(Obj));
        // Copy the object to the device (only has single scalar field to keep it simple)
        cudaMemcpy(h_d_obj[i], &(h_obj[i]), sizeof(Obj), cudaMemcpyHostToDevice);
    }

    /**************************************************/
    /* CREATE DEVICE ARRAY TO PASS POINTERS TO KERNEL */
    /**************************************************/

    // Create a pointer which will point to device memory
    Obj** d_d_obj = nullptr;
    // Allocate space for 3 pointers on device at above location
    cudaMalloc((void**)&d_d_obj, sizeof(Obj*) * 3);
    // Copy the pointers from the host memory to the device array
    cudaMemcpy(d_d_obj, h_d_obj, sizeof(Obj*) * 3, cudaMemcpyHostToDevice);


    /**********
     * After the above, VS2013 shows the memory pointed to by d_d_obj 
     * to be NULL <------- SEE QUESTION 2.
     **********/


    // Launch Kernel
    myKernel <<<1, 3>>>(d_d_obj);

    // Synchronise and pass back to host
    cudaDeviceSynchronize();
    for (int i = 0; i < 3; i++) {
        cudaMemcpy(&(h_obj[i]), h_d_obj[i], sizeof(Obj), cudaMemcpyDeviceToHost);     <--------- SEE QUESTION 3
    }

    // Write out
    for (int i = 0; i < 3; i++) {
        std::cout << h_obj[i]->scalar << std::endl;
    }

    return 0;
}

所以问题是：

1. 如果上面指示的行为SEE QUESTION 1指针分配主机内存，并且一旦我cudaMalloc在后续循环中使用分配设备内存，h_d_obj 指向的指针就会被设备地址覆盖，这是否意味着我已经为 3 分配了主机Obj*内存现在没有指针指向它吗？

2. 为什么cudaMemcpy我测试返回状态成功但明显没有正确复制地址？我期望两者的内存地址“数组”h_d_obj相同d_d_obj，因为它们应该Obj在设备地址空间中指向相同的地址。

3. 在这一行SEE QUESTION 3，假设我在问题 2 中是正确的。我还希望能够使用 或h_d_obj来从设备d_d_obj检索Obj对象，因为区别仅在于我是否取消引用主机指针来访问设备指针Obj或设备指针这两个我都可以在一个cudaMemcpy方法中完成，对吧？如果我使用写入的内容，复制会成功，但指针h_obj[0]已损坏，并且我无法写出数据。

4. 在该行SEE QUESTION 4，为什么我不能取消引用 anObj**来获取 aObj*然后使用->运算符来调用设备方法？编译器抱怨说它不是指向类类型的指针，但事实上它是一个类型，它Obj*告诉我它是。

Answer 1

首先，如果您提供完整的代码（包括类的定义）会很方便Obj。我根据对您的代码的检查和一些猜测提供了一个。

其次，您在这里的大部分困惑似乎是 C（或 C++）中的指针不太清晰。在主机和设备之间使用具有双指针结构 ( **) 的 CUDA API 需要清晰的理解和可视化正在发生的情况的能力。

如果上面的“参见问题1”指示的行为指针分配了主机内存，并且一旦我cudaMalloc在后续循环中使用了分配设备内存，所指向的指针就会被h_d_obj设备地址覆盖，这是否意味着我已经为指针分配了主机内存3 Obj* 现在没有指针指向它？

编号 h_d_obj是通过操作建立的（即赋予有意义的值）malloc。此后您所做的任何事情都不会修改 的值h_d_obj。

为什么当我测试返回的状态时，cudaMemcpy 成功，但显然没有正确复制地址？我期望两者的内存地址“数组”h_d_obj相同d_d_obj，因为它们应该指向设备地址空间中的相同 Obj。

到目前为止，我没有发现您的代码有任何问题。的值h_d_obj（之前）由 建立malloc，其数值是主机内存中的地址。的值d_d_obj由 建立cudaMalloc，其数值是设备内存中的地址。从数字上来说，我希望它们会有所不同。

在 SEE QUESTION 3 行，假设我在问题 2 中是正确的。我还希望能够使用 或h_d_obj从d_d_obj设备检索 Obj 对象，因为区别仅在于我是否取消引用主机指针来访问设备指向 Obj 的指针或设备指针，我都可以在 cudaMemcpy 方法中执行这两个操作，对吗？如果我使用写入的内容，复制会成功，但 h_obj[0] 处的指针已损坏，并且我无法写出数据。

不。您不能在主机代码中取消引用设备指针，即使它是cudaMemcpy. 作为操作中的源或目标，这是合法的cudaMemcpy：

h_d_obj[i]

这是不合法的：

d_d_obj[i]

原因是为了获得实际的目标地址，在第一种情况下我必须取消引用主机指针（即访问主机上的内存位置），但在第二种情况下必须取消引用设备指针。从主机代码中，我可以检索h_d_obj[i]. 我不允许尝试检索主机代码中的内容（以及主机代码的d_d_obj[i]参数操作）。的值可以用作主机代码的目标。 不能。cudaMemcpyd_d_objd_d_obj[i]

在 SEE QUESTION 4 行，为什么我不能取消引用 Obj** 来获取 Obj* 然后使用 -> 运算符调用设备方法？编译器抱怨说它不是指向类类型的指针，而它是 Obj* 的事实告诉我它是。

编译器对你咆哮，因为你不理解你正在使用的各种运算符（ *, ）之间的操作顺序。->如果添加括号来标识正确的顺序：

(*d_array_of_objs)->changeToFive(); 

那么编译器就不会反对这一点（尽管我的做法略有不同，如下所示）。

这是代码的修改版本Obj，添加了定义，对内核进行了轻微更改，以便独立线程在独立对象上工作，以及一些其他修复。您的代码大部分是正确的：

$ cat t1231.cu
#include <iostream>

class Obj{

  public:
  int scalar;
  __host__ __device__
  void changeToFive() {scalar = 5;}
};

// Kernel
__global__ void myKernel(Obj** d_array_of_objs)
{
    // Change the scalar of each object to 5
    // by dereferencing device array to get
    // appropriate object pointer.
    int idx = threadIdx.x+blockDim.x*blockIdx.x;
    // (*d_array_of_objs)->changeToFive();  //  <--------- SEE QUESTION 4 (add parenthesis)
    d_array_of_objs[idx]->changeToFive();
}

// Entry point
int main()
{

    /********************************/
    /* INITIALISE OBJ ARRAY ON HOST */
    /********************************/

    // Array of 3 pointers to Objs
    Obj* h_obj[3];
    for (int i = 0; i < 3; i++) {
        h_obj[i] = new Obj();       // Create
        h_obj[i]->scalar = i * 10;  // Initialise
    }

    // Write out
    for (int i = 0; i < 3; i++) {
        std::cout << h_obj[i]->scalar << std::endl;
    }


    /**************************************************/
    /* CREATE DEVICE VERSIONS AND STORE IN HOST ARRAY */
    /**************************************************/

    // Create host pointer to array-like storage of device pointers
    Obj** h_d_obj = (Obj**)malloc(sizeof(Obj*) * 3); //    <--------- SEE QUESTION 1
    for (int i = 0; i < 3; i++) {
        // Allocate space for an Obj and assign
        cudaMalloc((void**)&h_d_obj[i], sizeof(Obj));
        // Copy the object to the device (only has single scalar field to keep it simple)
        cudaMemcpy(h_d_obj[i], &(h_obj[i]), sizeof(Obj), cudaMemcpyHostToDevice);
    }

    /**************************************************/
    /* CREATE DEVICE ARRAY TO PASS POINTERS TO KERNEL */
    /**************************************************/

    // Create a pointer which will point to device memory
    Obj** d_d_obj = NULL;
    // Allocate space for 3 pointers on device at above location
    cudaMalloc((void**)&d_d_obj, sizeof(Obj*) * 3);
    // Copy the pointers from the host memory to the device array
    cudaMemcpy(d_d_obj, h_d_obj, sizeof(Obj*) * 3, cudaMemcpyHostToDevice);


    /**********
     * After the above, VS2013 shows the memory pointed to by d_d_obj
     * to be NULL <------- SEE QUESTION 2.
     **********/


    // Launch Kernel
    myKernel <<<1, 3>>>(d_d_obj);

    // Synchronise and pass back to host
    cudaDeviceSynchronize();
    for (int i = 0; i < 3; i++) {
        cudaMemcpy(h_obj[i], h_d_obj[i], sizeof(Obj), cudaMemcpyDeviceToHost);  //   <--------- SEE QUESTION 3  remove parenthesis
    }

    // Write out
    for (int i = 0; i < 3; i++) {
        std::cout << h_obj[i]->scalar << std::endl;
    }

    return 0;
}
$ nvcc -o t1231 t1231.cu
$ cuda-memcheck ./t1231
========= CUDA-MEMCHECK
0
10
20
5
5
5
========= ERROR SUMMARY: 0 errors
$

h_d_obj和的图表d_d_obj可能会有所帮助：

HOST                               |    DEVICE
h_d_obj-->(Obj *)-------------------------->Obj0<---(Obj *)<----|
          (Obj *)-------------------------->Obj1<---(Obj *)     |
          (Obj *)-------------------------->Obj2<---(Obj *)     |
                                   |                            |
d_d_obj---------------------------------------------------------|
HOST                               |    DEVICE

您可以访问上图左侧 (HOST)、主机代码或操作中的任何数量（位置）cudaMemcpy。您无法访问主机代码中右侧的任何数量（位置）。