Utiliser les données allouées dynamiquement dans le noyau CUDA sur l'hôte

Question

J'essaie de créer une classe de conteneur sur le périphérique qui gère de la mémoire. Cette mémoire est allouée dynamiquement et remplie lors de la construction de l'objet dans le noyau. Selon la documentation qui peut être faite avec un simple nouveau [] dans le noyau (en utilisant CUDA 8.0 avec la capacité de calcul 5.0 dans Visual Studio 2012). Ensuite, je veux accéder aux données à l'intérieur des conteneurs dans le code hôte (par exemple pour tester si toutes les valeurs sont correctes).

Une version minimale de la classe DeviceContainer ressemble à ceci:

class DeviceContainer 
{ 
public: 
    __device__ DeviceContainer(unsigned int size); 
    __host__ __device__ ~DeviceContainer(); 

    __host__ __device__ DeviceContainer(const DeviceContainer & other); 
    __host__ __device__ DeviceContainer & operator=(const DeviceContainer & other); 

    __host__ __device__ unsigned int getSize() const { return m_sizeData; } 
    __device__ int * getDataDevice() const { return mp_dev_data; } 
    __host__ int* getDataHost() const; 

private: 
    int * mp_dev_data; 
    unsigned int m_sizeData; 
}; 


__device__ DeviceContainer::DeviceContainer(unsigned int size) : 
     m_sizeData(size), mp_dev_data(nullptr) 
{ 
    mp_dev_data = new int[m_sizeData]; 

    for(unsigned int i = 0; i < m_sizeData; ++i) { 
     mp_dev_data[i] = i; 
    } 
} 


__host__ __device__ DeviceContainer::DeviceContainer(const DeviceContainer & other) : 
    m_sizeData(other.m_sizeData) 
{ 
#ifndef __CUDA_ARCH__ 
    cudaSafeCall(cudaMalloc((void**)&mp_dev_data, m_sizeData * sizeof(int))); 
    cudaSafeCall(cudaMemcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int), cudaMemcpyDeviceToDevice)); 
#else 
    mp_dev_data = new int[m_sizeData]; 
    memcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int)); 
#endif 
} 


__host__ __device__ DeviceContainer::~DeviceContainer() 
{ 
#ifndef __CUDA_ARCH__ 
    cudaSafeCall(cudaFree(mp_dev_data)); 
#else 
    delete[] mp_dev_data; 
#endif 
    mp_dev_data = nullptr; 
} 


__host__ __device__ DeviceContainer & DeviceContainer::operator=(const DeviceContainer & other) 
{ 
    m_sizeData = other.m_sizeData; 

#ifndef __CUDA_ARCH__ 
    cudaSafeCall(cudaMalloc((void**)&mp_dev_data, m_sizeData * sizeof(int))); 
    cudaSafeCall(cudaMemcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int), cudaMemcpyDeviceToDevice)); 
#else 
    mp_dev_data = new int[m_sizeData]; 
    memcpy(mp_dev_data, other.mp_dev_data, m_sizeData * sizeof(int)); 
#endif 

    return *this; 
} 


__host__ int* DeviceContainer::getDataHost() const 
{ 
    int * pDataHost = new int[m_sizeData]; 
    cudaSafeCall(cudaMemcpy(pDataHost, mp_dev_data, m_sizeData * sizeof(int), cudaMemcpyDeviceToHost)); 
    return pDataHost; 
} 

Il gère tout le tableau mp_dev_data. Le tableau est créé et rempli de valeurs consécutives pendant la construction, ce qui ne devrait être possible que sur l'appareil. (Notez qu'en réalité la taille des conteneurs peut être différente les uns des autres.)

Je pense que je dois fournir un constructeur de copie et un opérateur d'assignation puisque je ne connais pas d'autre moyen de remplir le tableau dans le noyau. (Voir la question n ° 3 ci-dessous.) Étant donné que la copie et la suppression peuvent également se produire sur l'hôte, __CUDA_ARCH__ est utilisé pour déterminer le chemin d'exécution que nous compilons. Sur l'hôte cudaMemcpy et cudaFree est utilisé, sur l'appareil, nous pouvons simplement utiliser memcpy et delete[].

Le noyau pour la création d'objets est assez simple:

__global__ void createContainer(DeviceContainer * pContainer, unsigned int numContainer, unsigned int containerSize) 
{ 
    unsigned int offset = blockIdx.x * blockDim.x + threadIdx.x; 

    if(offset < numContainer) 
    { 
     pContainer[offset] = DeviceContainer(containerSize); 
    } 
} 

Chaque fil dans une grille à une dimension qui est dans la gamme crée un seul objet conteneur.

La principale fonction alloue alors des tableaux pour le conteneur (90000 dans ce cas) sur l'appareil et l'hôte, appelle le noyau et tente d'utiliser les objets:

void main() 
{ 
    const unsigned int numContainer = 90000; 
    const unsigned int containerSize = 5; 

    DeviceContainer * pDevContainer; 
    cudaSafeCall(cudaMalloc((void**)&pDevContainer, numContainer * sizeof(DeviceContainer))); 

    dim3 blockSize(1024, 1, 1); 
    dim3 gridSize((numContainer + blockSize.x - 1)/blockSize.x , 1, 1); 

    createContainer<<<gridSize, blockSize>>>(pDevContainer, numContainer, containerSize); 
    cudaCheckError(); 

    DeviceContainer * pHostContainer = (DeviceContainer *)malloc(numContainer * sizeof(DeviceContainer)); 
    cudaSafeCall(cudaMemcpy(pHostContainer, pDevContainer, numContainer * sizeof(DeviceContainer), cudaMemcpyDeviceToHost)); 

    for(unsigned int i = 0; i < numContainer; ++i) 
    { 
     const DeviceContainer & dc = pHostContainer[i]; 

     int * pData = dc.getDataHost(); 
     for(unsigned int j = 0; j < dc.getSize(); ++j) 
     { 
     std::cout << pData[j]; 
     } 
     std::cout << std::endl; 
     delete[] pData; 
    } 

    free(pHostContainer); 
    cudaSafeCall(cudaFree(pDevContainer)); 
} 

Je dois utiliser malloc pour le tableau création sur l'hôte, puisque je ne veux pas avoir un constructeur par défaut pour le DeviceContainer. J'essaie d'accéder aux données à l'intérieur d'un conteneur via getDataHost() qui, en interne, appelle simplement cudaMemcpy.

cudaSafeCall et cudaCheckError sont des macros simples qui évaluent le cudaError retourné par la fonction ou interroger activement la dernière erreur. Par souci d'exhaustivité:

#define cudaSafeCall(error) __cudaSafeCall(error, __FILE__, __LINE__) 
#define cudaCheckError() __cudaCheckError(__FILE__, __LINE__) 

inline void __cudaSafeCall(cudaError error, const char *file, const int line) 
{ 
    if (error != cudaSuccess) 
    { 
     std::cerr << "cudaSafeCall() returned:" << std::endl; 
     std::cerr << "\tFile: " << file << ",\nLine: " << line << " - CudaError " << error << ":" << std::endl; 
     std::cerr << "\t" << cudaGetErrorString(error) << std::endl; 

     system("PAUSE"); 
     exit(-1); 
    } 
} 


inline void __cudaCheckError(const char *file, const int line) 
{ 
    cudaError error = cudaDeviceSynchronize(); 
    if (error != cudaSuccess) 
    { 
     std::cerr << "cudaCheckError() returned:" << std::endl; 
     std::cerr << "\tFile: " << file << ",\tLine: " << line << " - CudaError " << error << ":" << std::endl; 
     std::cerr << "\t" << cudaGetErrorString(error) << std::endl; 

     system("PAUSE"); 
     exit(-1); 
    } 
} 

J'ai 3 problèmes avec ce code:

1. Si elle est exécutée tel que présenté ici, je reçois un « échec du lancement non spécifié » du noyau. Le débogueur Nsight m'arrête sur la ligne mp_dev_data = new int[m_sizeData]; (soit dans le constructeur ou l'opérateur d'affectation) et signale plusieurs violations d'accès dans la mémoire globale. Le nombre de violations semble être aléatoire entre 4 et 11 et elles se produisent dans des threads non consécutifs mais toujours près de l'extrémité supérieure de la grille (blocs 85 et 86).

2. Si je réduis numContainer-10, le noyau se passe bien, cependant, le cudaMamcpy en getDataHost() échoue avec une erreur d'argument non valide - même si mp_dev_data est pas 0.(Je soupçonne que l'affectation est défectueuse et la mémoire a déjà été supprimée par un autre objet.)

3. Même si je voudrais savoir comment implémenter correctement le DeviceContainer avec une gestion de mémoire appropriée, dans mon cas, il serait également suffisant pour le rendre non-copiable et non-assignable. Cependant, je ne sais pas comment remplir correctement le conteneur-array dans le noyau. Peut-être quelque chose comme

   DeviceContainer dc(5); memcpy(&pContainer[offset], &dc, sizeof(DeviceContainer));

   Ce qui conduirait à des problèmes avec la suppression mp_dev_data dans le destructor. Je devrais gérer manuellement la suppression de mémoire qui se sent plutôt sale.

J'ai essayé aussi d'utiliser malloc et free dans le code du noyau au lieu de new et delete mais les résultats étaient les mêmes.

Je suis désolé de ne pas avoir pu formuler ma question plus rapidement. TL; DR: Comment implémenter une classe qui alloue dynamiquement de la mémoire dans un noyau et peut également être utilisée dans un code hôte? Comment puis-je initialiser un tableau dans un noyau avec des objets qui ne peuvent pas être copiés ou assignés?

Toute aide est appréciée. Je vous remercie.

Answer 1

Apparemment, la réponse est: Ce que j'essaie de faire est plus ou moins impossible. La mémoire allouée avec new ou malloc dans le noyau n'est pas placée dans la mémoire globale mais dans une mémoire de segment spéciale inaccessible depuis l'hôte. La seule option pour accéder à toute la mémoire sur l'hôte est d'allouer d'abord un tableau dans la mémoire globale qui est assez grand pour contenir tous les éléments sur le tas, puis écrire un noyau qui copie tous les éléments du tas vers la mémoire globale.

La violation d'accès est provoquée par la taille de segment limitée (qui peut être modifiée par