Utilisation d'un membre de la classe hôte pointant vers la mémoire du périphérique dans le code de périphérique

Question

Je souhaite avoir une instance d'une classe Container allouant une mémoire de périphérique et d'hôte lors de l'initialisation. Je veux utiliser la mémoire allouée dans le code de l'appareil, sans passer le pointeur réel (pour des raisons d'API).

Comment créer un pointeur global __device__ vers le membre pointant vers la mémoire de l'appareil? Je suis heureux d'utiliser la poussée si cela aide.

Voici un petit exemple:

#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

Il y a deux problèmes ici:

1. Vous ne pouvez pas inititalize statiquement une variable __device__ dans la façon dont vous essayez de (et la valeur essayez d'appliquer n'est pas correct non plus). L'API d'exécution CUDA contient une fonction d'initialisation des symboles de périphérique de portée globale. Utilisez cela à la place.
2. Votre déclaration de portée globale de stuff ne doit pas fonctionner pour un certain nombre de raisons subtiles abordées here (il s'agit d'un comportement techniquement indéfini). Déclarez-le au main champ à la place.

Mettre ces deux choses ensemble devrait conduire votre faire quelque chose comme ceci:

__device__ int* d_int; 

// ... 

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

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

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

    // ... 

Voici un exemple bien travaillé:

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