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

Question

我想要一个容器实例 class 在初始化时分配一些设备和主机内存。我想在设备代码中使用分配的内存，而不传递实际指针（出于 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

这里有两个问题：

您不能按照您尝试的方式静态初始化 __device__ 变量（并且您尝试应用的值也不正确）。 CUDA 运行时 API 包含一个用于初始化全局范围设备符号的函数。请改用它。
您的 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
$

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

Using host class member pointing to device memory in device code

cuda

thrust