Cuda对象副本

Question

我正在尝试将 CUDA 与对象一起使用，这是我放在一起进行尝试的一些测试代码，但我运行遇到了问题。当我对变量的设备版本做任何事情时，复制回主机失败并显示 "cuda Error Ilegal Address"，但如果我只是将代码复制到设备并返回它就可以了。如果我注释掉 printf... 行，它就可以了。

class A {
public:
    int s;
};

__device__ A *d_a;
__global__ void MethodA() {
    printf("%d\n", d_a->s);
}

int main() {
    A *a = new A();
    a->s = 10;

    cudaError e;
    e = cudaMalloc((void**)&d_a, sizeof(A));
    e = cudaMemcpy(d_a, a, sizeof(A), cudaMemcpyHostToDevice);
    MethodA << <1, 1 >> > ();
    e = cudaMemcpy(a, d_a, sizeof(A), cudaMemcpyDeviceToHost);
    std::cout << cudaGetErrorName(e) << std::endl;

    delete(a);
    std::getchar();
    return 0;
}

Answer 1

__device__ 变量的使用造成了困难。它旨在用于编译时已知的静态分配。

如果您使用普通的基于主机的指针，指向在运行时创建的动态分配（无论如何您都在做），然后将该基于主机的指针通过内核参数。

您的方法存在一些问题：

您使用了不正确的 API 来修改 __device__ 变量。我们不使用 cudaMemcpy。我们用cudaMemcpyToSymbol等
不允许在主机代码中获取设备实体的地址：
```
e = cudaMalloc((void**)&d_a, sizeof(A));
                       ^
```
cudaMalloc 期望将分配的指针值存储在 host 内存中，而不是设备内存中。它将指向设备内存中的一个位置，但它应该存储在一个主机变量中。

如果您想继续使用您的方法，请进行以下修改使其正确：

$ cat t89.cu
#include <iostream>
#include <stdio.h>

class A {
public:
    int s;
};

__device__ A *d_a;
__global__ void MethodA() {
    printf("%d\n", d_a->s);
}

int main() {
    A *a = new A();
    a->s = 10;
    A *temp_d_a;
    cudaMalloc((void**)&temp_d_a, sizeof(A));
    cudaMemcpy(temp_d_a, a, sizeof(A), cudaMemcpyHostToDevice);
    cudaMemcpyToSymbol(d_a, &temp_d_a, sizeof(A *));
    MethodA << <1, 1 >> > ();
    cudaMemcpy(a, temp_d_a, sizeof(A), cudaMemcpyDeviceToHost);
    std::cout << cudaGetErrorString(cudaGetLastError()) << std::endl;
    cudaFree(temp_d_a);
    delete(a);
    return 0;
}
$ nvcc t89.cu -o t89
$ cuda-memcheck ./t89
========= CUDA-MEMCHECK
10
no error
========= ERROR SUMMARY: 0 errors
$

编辑： 关于我之前的声明：

Your methodology would be simplified if you used an ordinary host-based pointer, pointing to a dynamic allocation created at runtime (which you are doing anyway), and then pass that host-based pointer to the device, via a kernel parameter.

并在下面的评论中被问及，这是一个展示该方法的有效示例：

$ cat t89.cu
#include <iostream>
#include <stdio.h>

class A {
public:
    int s;
};

__global__ void MethodA(A *a) {
    printf("%d\n", a->s);
}

int main() {
    A *a = new A();
    a->s = 10;
    A *d_a;  //  an ordinary host-based pointer
    cudaMalloc((void**)&d_a, sizeof(A)); //dynamic allocation created at runtime
    cudaMemcpy(d_a, a, sizeof(A), cudaMemcpyHostToDevice);
    MethodA << <1, 1 >> > (d_a);  // passed to kernel via parameter
    cudaMemcpy(a, d_a, sizeof(A), cudaMemcpyDeviceToHost);
    std::cout << cudaGetErrorString(cudaGetLastError()) << std::endl;
    cudaFree(d_a);
    delete(a);
    return 0;
}
$ nvcc -o t89 t89.cu
$ cuda-memcheck ./t89
========= CUDA-MEMCHECK
10
no error
========= ERROR SUMMARY: 0 errors
$

Cuda对象副本

Cuda object copy

cuda

object

memcpy