Comment allouer dynamicment des tableaux à l’intérieur d’un kernel?

J’ai besoin d’allouer dynamicment des tableaux à l’intérieur de la fonction du kernel. Comment puis-je faire ça?

Mon code est quelque chose comme ça:

__global__ func(float *grid_d,int n, int nn){ int i,j; float x[n],y[nn]; //Do some really cool and heavy computations here that takes hours. } 

Mais ça ne marchera pas. Si cela se trouvait dans le code de l’hôte, je pourrais utiliser malloc. cudaMalloc a besoin d’un pointeur sur l’hôte et d’un autre sur le périphérique. À l’intérieur de la fonction du kernel, je n’ai pas le pointeur hôte.

Donc qu’est ce que je devrais faire?

Si cela prend trop de temps (quelques secondes) pour allouer tous les tableaux (il me faut environ 4 de taille n et 5 de taille nn), ce ne sera pas un problème. Puisque le kernel fonctionnera probablement pendant au moins 20 minutes.

L’allocation dynamic de mémoire n’est prise en charge que par les capacités de calcul 2.x et ultérieures. Vous pouvez utiliser le nouveau mot-clé C ++ ou malloc dans le kernel pour que votre exemple devienne:

 __global__ func(float *grid_d,int n, int nn){ int i,j; float *x = new float[n], *y = new float[nn]; } 

Cela alloue de la mémoire sur un segment d’exécution de la mémoire locale qui a la durée de vie du contexte. Assurez-vous donc de libérer la mémoire une fois le kernel exécuté si votre intention n’est pas de l’utiliser à nouveau. Vous devez également noter qu’il est impossible d’accéder directement à la mémoire de cudaMemcpy d’ cudaMemcpy à partir des API hôtes. Par conséquent, vous ne pouvez pas transmettre un pointeur alloué dans un kernel en tant qu’argument à cudaMemcpy , par exemple.

@talonmies a répondu à votre question sur l’allocation dynamic de mémoire dans un kernel. Ceci est conçu comme une réponse supplémentaire, abordant les performances de __device__ malloc() et une alternative à envisager.

L’allocation dynamic de mémoire dans le kernel peut être tentante, car elle permet au code GPU de ressembler davantage à du code CPU. Mais cela peut sérieusement affecter les performances. J’ai écrit un test autonome et je l’ai inclus ci-dessous. Le test lance quelque 2,6 millions de threads. Chaque thread remplit 16 entiers de la mémoire globale avec des valeurs dérivées de l’index du thread, puis résume les valeurs et renvoie la sum.

Le test met en œuvre deux approches. La première approche utilise __device__ malloc() et la seconde utilise la mémoire allouée avant l’exécution du kernel.

Sur mon périphérique 2.0, le kernel s’exécute en 1500 ms avec __device__ malloc() et 27 ms avec de la mémoire préallouée. En d’autres termes, l’exécution du test prend 56 fois plus longtemps lorsque la mémoire est allouée de manière dynamic dans le kernel. Le temps inclut la boucle externe cudaMalloc() / cudaFree() , qui ne fait pas partie du kernel. Si le même kernel est lancé plusieurs fois avec le même nombre de threads, comme c’est souvent le cas, le coût de cudaMalloc() / cudaFree() est amorti sur tous les lancements du kernel. Cela porte la différence encore plus haut, à environ 60x.

En spéculant, je pense que la baisse des performances est en partie due à la sérialisation implicite. Le processeur graphique doit probablement sérialiser tous les appels simultanés à __device__ malloc() afin de fournir des blocs de mémoire distincts à chaque appelant.

La version qui n’utilise pas __device__ malloc() alloue toute la mémoire du processeur graphique avant d’exécuter le kernel. Un pointeur sur la mémoire est transmis au kernel. Chaque thread calcule un index dans la mémoire allouée précédemment au lieu d’utiliser un __device__ malloc() .

Le problème potentiel avec l’allocation de mémoire à l’avance est que, si seuls certains threads doivent allouer de la mémoire et qu’on ne sait pas quels sont ces threads, il sera nécessaire d’allouer de la mémoire pour tous les threads. S’il n’y a pas assez de mémoire pour cela, il serait peut-être plus efficace de réduire le nombre de threads par appel du kernel que d’utiliser __device__ malloc() . D’autres solutions de contournement finiraient probablement par réimplémenter ce que __device__ malloc() fait en arrière-plan et entraîneraient un __device__ malloc() similaire sur les performances.

Testez les performances de __device__ malloc() :

 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include  const int N_ITEMS(16); #define USE_DYNAMIC_MALLOC __global__ void test_malloc(int* totals) { int tx(blockIdx.x * blockDim.x + threadIdx.x); int* s(new int[N_ITEMS]); for (int i(0); i < N_ITEMS; ++i) { s[i] = tx * i; } int total(0); for (int i(0); i < N_ITEMS; ++i) { total += s[i]; } totals[tx] = total; delete[] s; } __global__ void test_malloc_2(int* items, int* totals) { int tx(blockIdx.x * blockDim.x + threadIdx.x); int* s(items + tx * N_ITEMS); for (int i(0); i < N_ITEMS; ++i) { s[i] = tx * i; } int total(0); for (int i(0); i < N_ITEMS; ++i) { total += s[i]; } totals[tx] = total; } int main() { cudaError_t cuda_status; cudaSetDevice(0); int blocks_per_launch(1024 * 10); int threads_per_block(256); int threads_per_launch(blocks_per_launch * threads_per_block); int* totals_d; cudaMalloc((void**)&totals_d, threads_per_launch * sizeof(int)); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaDeviceSynchronize(); cudaEventRecord(start, 0); #ifdef USE_DYNAMIC_MALLOC cudaDeviceSetLimit(cudaLimitMallocHeapSize, threads_per_launch * N_ITEMS * sizeof(int)); test_malloc<<>>(totals_d); #else int* items_d; cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS); test_malloc_2<<>>(items_d, totals_d); cudaFree(items_d); #endif cuda_status = cudaDeviceSynchronize(); if (cuda_status != cudaSuccess) { printf("Error: %d\n", cuda_status); exit(1); } cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop); printf("Elapsed: %f\n", elapsedTime); int* totals_h(new int[threads_per_launch]); cuda_status = cudaMemcpy(totals_h, totals_d, threads_per_launch * sizeof(int), cudaMemcpyDeviceToHost); if (cuda_status != cudaSuccess) { printf("Error: %d\n", cuda_status); exit(1); } for (int i(0); i < 10; ++i) { printf("%d ", totals_h[i]); } printf("\n"); cudaFree(totals_d); delete[] totals_h; return cuda_status; } 

Sortie:

 C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe Elapsed: 27.311169 0 120 240 360 480 600 720 840 960 1080 C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe Elapsed: 1516.711914 0 120 240 360 480 600 720 840 960 1080 

Si les valeurs de n et nn étaient connues avant l’appel du kernel, pourquoi ne pas cudaMalloc placer la mémoire côté hôte et transmettre le pointeur de mémoire de périphérique au kernel?

Une expérience basée sur les concepts du billet de @ rogerdahl. Hypothèses:

  • 4 Mo de mémoire allouée en morceaux de 64 Go.
  • 1 bloc GPU et 32 ​​fils de chaîne dans ce bloc
  • Courir sur un P100

Les appels gratuits malloc + locaux au GPU semblaient beaucoup plus rapides que les cudaMalloc + cudaFree . La sortie du programme:

 Starting timer for cuda malloc timer Stopping timer for cuda malloc timer timer for cuda malloc timer took 1.169631s Starting timer for device malloc timer Stopping timer for device malloc timer timer for device malloc timer took 0.029794s 

Je laisse le code pour timer.h et timer.cpp , mais voici le code du test lui-même:

 #include "cuda_runtime.h" #include  #include  #include "timer.h" static void CheckCudaErrorAux (const char *, unsigned, const char *, cudaError_t); #define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value) const int BLOCK_COUNT = 1; const int THREADS_PER_BLOCK = 32; const int ITERATIONS = 1 << 12; const int ITERATIONS_PER_BLOCKTHREAD = ITERATIONS / (BLOCK_COUNT * THREADS_PER_BLOCK); const int ARRAY_SIZE = 64; void CheckCudaErrorAux (const char *file, unsigned line, const char *statement, cudaError_t err) { if (err == cudaSuccess) return; std::cerr << statement<<" returned " << cudaGetErrorString(err) << "("<>>(); CUDA_CHECK_RETURN(cudaDeviceSynchronize()); device_malloc_timer.stop_and_report(); } 

Si vous trouvez des erreurs, merci de les inclure dans les commentaires, et je vais essayer de les corriger.

Et je les ai de nouveau courus avec un plus gros

 const int BLOCK_COUNT = 56; const int THREADS_PER_BLOCK = 1024; const int ITERATIONS = 1 << 18; const int ITERATIONS_PER_BLOCKTHREAD = ITERATIONS / (BLOCK_COUNT * THREADS_PER_BLOCK); const int ARRAY_SIZE = 1024; 

Et cudaMalloc était encore beaucoup plus lent:

 Starting timer for cuda malloc timer Stopping timer for cuda malloc timer timer for cuda malloc timer took 74.878016s Starting timer for device malloc timer Stopping timer for device malloc timer timer for device malloc timer took 0.167331s 

Je suis nouveau ici. Je ne sais pas comment vous contacter, ragerdl.

Peut-être devriez-vous tester

 cudaMalloc(&foo,sizeof(int) * ARRAY_SIZE * ITERATIONS); cudaFree(foo); 

au lieu

 for (int i = 0; i < ITERATIONS; ++ i) { if (i == 1) cuda_malloc_timer.start(); // let it warm up one cycle int * foo; cudaMalloc(&foo, sizeof(int) * ARRAY_SIZE); cudaFree(foo); }