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__ 变量会造成困难。它旨在用于静态分配，在编译时已知。

如果您使用普通的基于主机的指针，指向在运行时创建的动态分配（无论如何您都在这样做），然后通过内核参数将该基于主机的指针传递给设备，您的方法将得到简化.

您的方法存在一些问题：

1. 您正在使用不正确的 API 来修改 __device__ 变量。我们不使用cudaMemcpy。我们使用cudaMemcpyToSymbol等。

2. 您不能在主机代码中获取设备实体的地址：

   e = cudaMalloc((void**)&d_a, sizeof(A));
                          ^
   

   cudaMalloc 期望将分配的指针值存储在 主机 内存中，而不是设备内存中。它将指向设备内存中的一个位置，但它应该存储在一个主机变量中。

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

$ 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
$

编辑：关于我之前的陈述：

如果您使用普通的基于主机的指针，指向在运行时创建的动态分配（无论如何您都在这样做），然后通过内核参数将该基于主机的指针传递给设备，您的方法将得到简化.

并在下面的 cmets 中询问，这是一个展示该方法的工作示例：

$ 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
$