在设备代码中使用指向设备内存的主机类成员

Question

我希望有一个Container类的实例在初始化时分配一些设备和主机内存。我想在设备代码中使用分配的内存，而不传递实际的指针（API原因）。

如何创建全局__device__指向指向设备内存的成员的指针？如果有帮助，我很乐意使用推力。

这里是一个小例子：

#include <iostream> 


struct Container { 
    int *h_int = (int*)malloc(4*sizeof(int)); 
    int *d_int; 
    Container() { 
     h_int[0] = 6; h_int[1] = 6; h_int[2] = 6; h_int[3] = 6; 
     cudaMalloc(&d_int, 4*sizeof(int)); 
     memcpyHostToDevice(); 
    } 
    void memcpyHostToDevice() { 
     cudaMemcpy(d_int, h_int, 4*sizeof(int), cudaMemcpyHostToDevice); 
    } 
    void memcpyDeviceToHost() { 
     cudaMemcpy(h_int, d_int, 4*sizeof(int), cudaMemcpyDeviceToHost); 
    } 
}; 

Container stuff; 


__device__ auto d_int = &stuff.d_int; // How do I get that right? 


__global__ void edit() { // To keep the API simple I do not want to pass the pointer 
    auto i = blockIdx.x*blockDim.x + threadIdx.x; 
    d_int[i] = 1 + 2*(i > 0) + 4*(i > 2); 
} 


int main(int argc, char const *argv[]) { 
    edit<<<4, 1>>>(); 
    stuff.memcpyDeviceToHost(); 
    std::cout << stuff.h_int[0] << stuff.h_int[1] << stuff.h_int[2] << stuff.h_int[3] << "\n"; 
    return 0; 
} 

Answer 1

有两个问题在这里：

1. 不能静态inititalize在您尝试的方式__device__变量（和你的价值试图申请也是不正确的）。 CUDA运行时API包含一个用于初始化全局范围设备符号的函数。改为使用它。
2. 您的全球范围声明stuff不应用于讨论here（这在技术上未定义的行为）讨论的一些微妙的原因。改为在main范围内声明它。

把这两件事一起应该带领你做这样的事情，而不是：

__device__ int* d_int; 

// ... 

int main(int argc, char const *argv[]) { 

    Container stuff; 
    cudaMemcpyToSymbol(d_int, &stuff.dint, sizeof(int*)); 

    edit<<<4, 1>>>(); 

    // ... 

这里是一个完全样例：

$ cat t1199.cu 
#include <iostream> 


struct Container { 
    int *h_int = (int*)malloc(4*sizeof(int)); 
    int *d_int; 
    Container() { 
     h_int[0] = 6; h_int[1] = 6; h_int[2] = 6; h_int[3] = 6; 
     cudaMalloc(&d_int, 4*sizeof(int)); 
     memcpyHostToDevice(); 
    } 
    void memcpyHostToDevice() { 
     cudaMemcpy(d_int, h_int, 4*sizeof(int), cudaMemcpyHostToDevice); 
    } 
    void memcpyDeviceToHost() { 
     cudaMemcpy(h_int, d_int, 4*sizeof(int), cudaMemcpyDeviceToHost); 
    } 
}; 

//Container stuff; 


__device__ int *d_int; // = &stuff.d_int; // How do I get that right? 


__global__ void edit() { // To keep the API simple I do not want to pass the pointer 
    auto i = blockIdx.x*blockDim.x + threadIdx.x; 
    d_int[i] = 1 + 2*(i > 0) + 4*(i > 2); 
} 


int main(int argc, char const *argv[]) { 
    Container stuff; 
    cudaMemcpyToSymbol(d_int, &stuff.d_int, sizeof(int *)); 
    edit<<<4, 1>>>(); 
    stuff.memcpyDeviceToHost(); 
    std::cout << stuff.h_int[0] << stuff.h_int[1] << stuff.h_int[2] << stuff.h_int[3] << "\n"; 
    return 0; 
} 
$ nvcc -std=c++11 -o t1199 t1199.cu 
$ cuda-memcheck ./t1199 
========= CUDA-MEMCHECK 
1337 
========= ERROR SUMMARY: 0 errors 
$