Comment allouer dynamiquement des tableaux à l'intérieur d'un noyau?

@talonmies répondu à votre question sur la façon d'allouer dynamiquement de la mémoire au sein d'un noyau. C'est prévu pour compléter la réponse, traitant de la performance de __device__ malloc() et alternative, vous pourriez envisager.

L'allocation dynamique de la mémoire dans le noyau peut être tentant, car il permet de code GPU à regarder de plus comme le code de CPU. Mais il peut sérieusement affecter les performances. J'ai écrit un autonome de test, et nous avons inclus ci-dessous. Le test des lancements quelque 2,6 millions de threads. Chaque thread remplit 16 entiers de la mémoire globale avec certaines valeurs issues du fil de l'index, puis résume les valeurs et renvoie la somme.

Le test met en œuvre deux approches. La première approche utilise __device__ malloc() et la deuxième approche utilise de la mémoire qui est allouée avant que le noyau s'exécute.

Sur mon 2.0 appareil, le noyau s'exécute dans 1500ms lors de l'utilisation de __device__ malloc() et 27 ms est observée lors de l'utilisation de pré-alloué de la mémoire. En d'autres termes, le test prend 56x plus à exécuter lorsque la mémoire est allouée dynamiquement à l'intérieur du noyau. Le temps comprend la boucle externe cudaMalloc() /cudaFree(), qui ne fait pas partie du noyau. Si le même noyau est lancé plusieurs fois avec le même nombre de fils, comme c'est souvent le cas, le coût de la cudaMalloc() /cudaFree() est amorti sur tous le noyau lance. Qu'apporte la différence encore plus élevé, à environ 60x.

De spéculer, je pense que le gain de performance est en partie causé par implicite de sérialisation. Le GPU doit probablement sérialiser tous les appels simultanés à __device__ malloc() afin de fournir des morceaux de mémoire pour chaque appelant.

La version qui n'utilise pas __device__ malloc() alloue tous les GPU de la mémoire avant de lancer le noyau. Un pointeur vers la mémoire est transmis au noyau. Chaque thread calcule un indice dans le précédemment alloué de la mémoire au lieu d'utiliser un __device__ malloc().

Le problème potentiel avec l'allocation de mémoire à l'avant, c'est que, si seuls certains threads ont besoin d'allouer de la mémoire, et on ne sait pas qui threads ceux qui sont, il sera nécessaire d'allouer de la mémoire pour tous les threads. Si il n'y a pas assez de mémoire pour que, il pourrait être plus efficace de réduire le nombre de threads par noyau appel, puis à l'aide de __device__ malloc(). D'autres solutions de contournement serait probablement jusqu'à la fin de réimplanter ce __device__ malloc() est fait en arrière-plan, et serait de voir une performance similaire a frappé.

Tester les performances de __device__ malloc():

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
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<<<blocks_per_launch, threads_per_block>>>(totals_d);
#else
int* items_d;
cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS);
test_malloc_2<<<blocks_per_launch, threads_per_block>>>(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;
}

De 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

Mené une expérience basée sur les concepts de la @rogerdahl post. Hypothèses:

• De 4 mo de mémoire allouée dans 64B morceaux.
• 1 GPU bloc et 32 fils de chaîne dans le canton
• Exécuter sur un P100

Le malloc+appels gratuits local pour le GPU semble être beaucoup plus rapide que la cudaMalloc + cudaFree appels. Le programme de la sortie:

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 pars le code pour timer.h et timer.cpp, mais voici le code pour le test lui-même:

#include "cuda_runtime.h"
#include <stdio.h>
#include <thrust/system/cuda/error.h>
#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) << "("<<err<< ") at "<<file<<":"<<line << std::endl;
exit (1);
}
__global__ void mallocai() {
for (int i = 0; i < ITERATIONS_PER_BLOCKTHREAD; ++i) {
int * foo;
foo = (int *) malloc(sizeof(int) * ARRAY_SIZE);
free(foo);
}
}
int main() {
Timer cuda_malloc_timer("cuda malloc timer");
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);
}
cuda_malloc_timer.stop_and_report();
CUDA_CHECK_RETURN(cudaDeviceSynchronize());
Timer device_malloc_timer("device malloc timer");
device_malloc_timer.start();
mallocai<<<BLOCK_COUNT, THREADS_PER_BLOCK>>>();
CUDA_CHECK_RETURN(cudaDeviceSynchronize());
device_malloc_timer.stop_and_report();
}

Si vous trouvez des erreurs, veuillez lmk dans les commentaires, et je vais essayer de les réparer.

Et j'ai couru de nouveau avec un plus grand tout:

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 plus lent par beaucoup:

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