Pourquoi n'y a-t-il aucune activité sur le GPU entre le tri successif et les commandes de réduction?

Question

S'il vous plaît se référer aux deux instantanés ci-dessous montrant une Nvidia visuelle session de Profiler de mon code CUDA:

Aperçu de la session nvprof montrant poussée :: sort et poussée :: réduire calendrier d'exécution d'appel

Mis en évidence le genre et réduire les appels à montrer les temps pris et l'écart entre leur exécution

Vous pouvez voir un écart d'environ 7 0 nous entre les deux appels thrust::sort(), puis il ya un grand écart entre le premier thrust::reduce() et le deuxième thrust::sort() appels. En tout il y a environ 300 us de tels trous visibles dans l'instantané. Je crois que ceux-ci sont 'ralenti' fois, peut-être introduit par la bibliothèque de poussée. Quoi qu'il en soit, je n'ai trouvé aucune discussion pertinente, ni aucune documentation à ce sujet par Nvidia. Quelqu'un peut-il expliquer s'il vous plaît pourquoi ai-je si apparente 'ralenti' fois? Combiné, ces temps représentent 40% du temps d'exécution de mon application, donc c'est une grande préoccupation pour moi!

En outre, j'ai mesuré que les intervalles entre les appels aux noyaux cuda successifs que j'ai écrit est d'environ 3 us!

J'ai écrit un exemple de code cuda pour poster ici:

void profileThrustSortAndReduce(const int ARR_SIZE) { 
    // for thrust::reduce on first 10% of the sorted array 
    const int ARR_SIZE_BY_10 = ARR_SIZE/10; 

    // generate host random arrays of float 
    float* h_arr1;   cudaMallocHost((void **)&h_arr1, ARR_SIZE * sizeof(float)); 
    float* h_arr2;   cudaMallocHost((void **)&h_arr2, ARR_SIZE * sizeof(float)); 
    for (int i = 0; i < ARR_SIZE; i++) { 
     h_arr1[i] = static_cast <float> (rand())/static_cast <float> (RAND_MAX)* 1000.0f; 
     h_arr2[i] = static_cast <float> (rand())/static_cast <float> (RAND_MAX)* 1000.0f; 
    } 

    // device arrays populated 
    float* d_arr1;   cudaMalloc((void **)&d_arr1, ARR_SIZE * sizeof(float)); 
    float* d_arr2;   cudaMalloc((void **)&d_arr2, ARR_SIZE * sizeof(float)); 
    cudaMemcpy(d_arr1, h_arr1, ARR_SIZE * sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_arr2, h_arr2, ARR_SIZE * sizeof(float), cudaMemcpyHostToDevice); 

    // start cuda profiler 
    cudaProfilerStart(); 

    // sort the two device arrays 
    thrust::sort(thrust::device, d_arr1, d_arr1 + ARR_SIZE); 
    thrust::sort(thrust::device, d_arr2, d_arr2 + ARR_SIZE); 

    // mean of 100 percentiles of device array 
    float arr1_red_100pc_mean = thrust::reduce(thrust::device, d_arr1, d_arr1 + ARR_SIZE)/ARR_SIZE; 
    // mean of smallest 10 percentiles of device array 
    float arr1_red_10pc_mean = thrust::reduce(thrust::device, d_arr1, d_arr1 + ARR_SIZE_BY_10)/ARR_SIZE_BY_10; 

    // mean of 100 percentiles of device array 
    float arr2_red_100pc_mean = thrust::reduce(thrust::device, d_arr2, d_arr2 + ARR_SIZE)/ARR_SIZE; 
    // mean of smallest 10 percentiles of device array 
    float arr2_red_10pc_mean = thrust::reduce(thrust::device, d_arr2, d_arr2 + ARR_SIZE_BY_10)/ARR_SIZE_BY_10; 

    // stop cuda profiler 
    cudaProfilerStop(); 
} 

Aperçu de la session nvprof de cette fonction exemple

Answer 1

Les lacunes sont principalement causées par cudaMalloc opérations. thrust::sort et vraisemblablement thrust::reduce allouer (et gratuit) le stockage temporaire associé à leur activité.

Vous avez supprimé cette partie de la timeline des 2 premières images que vous avez collées dans votre question, mais immédiatement au-dessus de la partie de la timeline que vous avez montrée dans votre 3ème image, vous trouverez cudaMalloc opérations dans le " runtime API "ligne de profileur. Ces opérations cudaMalloc (et cudaFree) prennent beaucoup de temps et se synchronisent. Pour contourner ce problème, le conseil typique est d'utiliser un thrust custom allocator (aussi here). Ce faisant, vous pouvez allouer une fois pour les tailles nécessaires, au début de votre programme, et ne pas avoir à engager l'allocation/freehead chaque fois que vous faites un appel de poussée. Vous pouvez également explorer cub, qui a déjà les étapes d'allocation et de traitement séparées pour vous.