在設備代碼中使用指向設備內存的主機類成員

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 
$