Host-Klasse Mitglied Verwendung in Gerätecode auf dem Gerätespeicher zeigt

Question

Ich mag eine Instanz einer Container-Klasse habe einige Geräte und Host-Speicher bei der Initialisierung zugewiesen wird. Ich möchte den zugewiesenen Speicher im Gerätecode verwenden, ohne den eigentlichen Zeiger zu übergeben (aus API-Gründen).

Wie kann ich einen globalen __device__ Zeiger auf das Element erstellen, um den Gerätespeicher zeigt? Ich bin glücklich, Schub zu verwenden, wenn das hilft.

Hier ist ein kleines Beispiel:

#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

Es gibt zwei Probleme:

1. Sie nicht statisch eine __device__ Variable in der Art und Weise Sie versuchen (und den Wert inititalize können Sie versuchen zu bewerben ist auch nicht korrekt). Die CUDA-Laufzeit-API enthält eine Funktion zum Initialisieren globaler Gerätesymbole. Benutze das stattdessen.
2. Ihre globale Reichweite Erklärung stuff sollte auch nicht für eine Reihe von subtilen Gründen arbeitet here diskutiert (es ist technisch nicht definiertes Verhalten). Erklären Sie es stattdessen unter main.

Setzt man diese beiden Dinge zusammen sollten Ihre führen so etwas wie dies zu tun, statt:

__device__ int* d_int; 

// ... 

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

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

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

    // ... 

Hier ist ein voll gearbeitet Beispiel:

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