Tableaux dynamiques CUDA globaux (comme en C) alloués à la mémoire de l'appareil

Question

Donc, je suis en train d’essayer d’écrire du code qui utilise l’architecture CUDA de Nvidia. J'ai remarqué que la copie vers et depuis l'appareil nuisait vraiment à mes performances globales. J'essaie donc de transférer une grande quantité de données sur l'appareil.

Comme ces données sont utilisées dans de nombreuses fonctions, je voudrais qu’elles soient globales. Oui, je peux faire passer des pointeurs, mais j'aimerais vraiment savoir comment travailler avec des globals dans ce cas.

J'ai donc des fonctions de périphérique qui veulent accéder à un tableau alloué par périphérique.

Idéalement, je pourrais faire quelque chose comme:

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

Cependant, je n'ai pas trouvé comment créer un tableau dynamique. J'ai trouvé un moyen de contourner le problème en déclarant le tableau comme suit:

__device__ float global_data[REALLY_LARGE_NUMBER];

Bien que cela ne nécessite pas d'appels cudaMalloc, je préférerais l'approche d'allocation dynamique.

Answer 1

Quelque chose comme ça devrait probablement fonctionner.

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)


__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}


int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

Faites un tourbillon.

Answer 2

Passez un peu de temps à vous concentrer sur la documentation fournie par NVIDIA.

À partir du Guide de programmation:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

Voilà un exemple simple de la façon d'allouer de la mémoire. Maintenant, dans vos noyaux, vous devriez accepter un pointeur sur un float comme ceci:

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

Alors maintenant, vous pouvez les invoquer comme suit:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

  

Comme ces données sont utilisées dans de nombreux   fonctions, je voudrais qu'il soit   global.

Il y a peu de bonnes raisons d'utiliser des globals. Ce n'en est certainement pas un. Je vais laisser cela comme un exercice pour développer cet exemple afin d'inclure le déplacement de "& dev; Ptr". à une portée globale.

EDIT:

Ok, le problème fondamental est le suivant: vos noyaux ne peuvent accéder qu'à la mémoire de l'appareil et les seuls pointeurs à portée globale qu'ils peuvent utiliser sont ceux utilisant un processeur graphique. Lorsque vous appelez un noyau à partir de votre CPU, il se produit en coulisse que les pointeurs et les primitives sont copiés dans des registres GPU et / ou dans la mémoire partagée avant que le noyau ne soit exécuté.

Le point le plus proche que je puisse suggérer est le suivant: utilisez cudaMemcpyToSymbol () pour atteindre vos objectifs. Mais, à l’arrière-plan, considérez qu’une approche différente pourrait être la bonne chose.

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}


int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

N'oubliez pas '--host-compilation = c ++' pour cet exemple.

Answer 3

Je suis allé de l'avant et j'ai essayé la solution consistant à allouer un pointeur temporaire et à le transmettre à une fonction globale simple similaire à kernel1.

La bonne nouvelle est que cela fonctionne:)

Cependant, je pense que cela confond le compilateur car je reçois à présent le message "Avis: Je ne peux pas dire à quel point pointe le pointeur, en supposant un espace mémoire global". chaque fois que j'essaie d'accéder aux données globales. Heureusement, cette hypothèse est correcte, mais les avertissements sont agaçants.

Quoi qu’il en soit, aux fins du compte rendu, j’ai examiné de nombreux exemples et j’ai parcouru les exercices nvidia où il s’agit de faire en sorte que la sortie dise "Correct!". Cependant, je n'ai pas examiné tous d'entre eux. Si quelqu'un connaît un exemple de SDK dans lequel il alloue dynamiquement la mémoire d'un périphérique global, j'aimerais tout de même savoir.

Answer 4

Heu, c’est précisément le problème de déplacement de devPtr vers une portée globale qui était mon problème.

J'ai une implémentation qui fait exactement cela, avec les deux noyaux ayant un pointeur sur les données passées. Je ne veux pas explicitement transmettre ces pointeurs.

J'ai lu assez attentivement la documentation et consulté les forums nvidia (et Google a cherché pendant environ une heure), mais je n'ai pas trouvé d'implémentation d'un tableau de périphériques dynamique global qui fonctionne (j'ai essayé plusieurs compiler puis échouer de manière nouvelle et intéressante).

Answer 5

Découvrez les exemples fournis avec le SDK. Nombre de ces exemples de projets sont une manière décente d'apprendre par l'exemple.

Answer 6

  

Comme ces données sont utilisées dans de nombreuses fonctions, j'aimerais qu'elles soient globales.

-

  

Il y a peu de bonnes raisons d'utiliser des globals. Ce n'en est certainement pas un. Je vais le laisser comme un   Exercice pour développer cet exemple pour inclure le déplacement " devPtr " à une portée globale.

Que se passe-t-il si le noyau fonctionne sur une grande structure const constituée de tableaux? L'utilisation de la mémoire dite constante n'est pas une option, car sa taille est très limitée .. vous devez donc la mettre en mémoire globale ..?