Comment puis-je accéder aux ID de flux numériques vus dans nvprof, en utilisant un cudaStream_t?

Question

Dans nvprof je peux voir les ID de flux pour chaque flux d'exécution cuda que j'utilise (0, 13, 15, etc.)

Étant donné une variable de flux, j'aimerais pouvoir imprimer l'identifiant de flux . Actuellement, je ne trouve aucune API pour ce faire et la conversion du cudaStream_t en un int ou en uint ne donne pas un ID raisonnable. sizeof() dit cudaStream_t est de 8 octets.

Answer 1

En bref: Je ne connais pas de méthode pour accéder directement à ces ID, mais vous pouvez donner des noms explicites aux flux à des fins de profilage. Type

---

cudaStream_t est opaque "poignée de ressources". Un handle de ressource est quelque chose comme un pointeur; il va donc de soi que l'ID de flux n'est pas contenu dans le pointeur (handle) lui-même, mais en quelque sorte dans ce à quoi il se réfère. Puisqu'il est opaque (aucune définition de ce à quoi il pointe, fourni par CUDA) et comme vous le signalez il n'y a aucune API directe pour cela, je ne pense pas que vous trouviez une méthode pour extraire l'ID de flux à partir d'un cudaStream_t à l'exécution.

Pour ces affirmations que cudaStream_t est une poignée de ressources et qu'il est opaque, reportez-vous au fichier d'en-tête CUDA driver_types.h

Cependant, the NV Tools Extension API vous donne la possibilité de « nom » un flux particulier (ou d'autres ressources). Cela vous permet d'associer un flux particulier dans le code source à un nom particulier dans le profileur.

Voici un exemple trivial travaillé:

$ cat t138.cu 
#include <stdio.h> 
#include <nvToolsExtCudaRt.h> 
const long tdel = 1000000000ULL; 
__global__ void tkernel(){ 

    long st = clock64(); 
    while (clock64() < st+tdel); 
} 

int main(){ 

    cudaStream_t s1, s2, s3, s4; 
    cudaStreamCreate(&s1); 
    cudaStreamCreate(&s2); 
    cudaStreamCreate(&s3); 
    cudaStreamCreate(&s4); 
#ifdef USE_S_NAMES 
    nvtxNameCudaStreamA(s1, "stream 1"); 
    nvtxNameCudaStreamA(s2, "stream 2"); 
    nvtxNameCudaStreamA(s3, "stream 3"); 
    nvtxNameCudaStreamA(s4, "stream 4"); 
#endif 
    tkernel<<<1,1,0,s1>>>(); 
    tkernel<<<1,1,0,s2>>>(); 
    tkernel<<<1,1,0,s3>>>(); 
    tkernel<<<1,1,0,s4>>>(); 

    cudaDeviceSynchronize(); 
} 

$ nvcc -arch=sm_61 -o t138 t138.cu -lnvToolsExt 
$ nvprof --print-gpu-trace ./t138 
==28720== NVPROF is profiling process 28720, command: ./t138 
==28720== Profiling application: ./t138 
==28720== Profiling result: 
    Start Duration   Grid Size  Block Size  Regs* SSMem* DSMem*  Size Throughput   Device Context Stream Name 
464.80ms 622.06ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1  13 tkernel(void) [393] 
464.81ms 621.69ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1  14 tkernel(void) [395] 
464.82ms 623.30ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1  15 tkernel(void) [397] 
464.82ms 622.69ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1  16 tkernel(void) [399] 

Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows. 
SSMem: Static shared memory allocated per CUDA block. 
DSMem: Dynamic shared memory allocated per CUDA block. 
$ nvcc -arch=sm_61 -o t138 t138.cu -lnvToolsExt -DUSE_S_NAMES 
$ nvprof --print-gpu-trace ./t138 
==28799== NVPROF is profiling process 28799, command: ./t138 
==28799== Profiling application: ./t138 
==28799== Profiling result: 
    Start Duration   Grid Size  Block Size  Regs* SSMem* DSMem*  Size Throughput   Device Context Stream Name 
457.98ms 544.07ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1 stream 1 tkernel(void) [393] 
457.99ms 544.31ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1 stream 2 tkernel(void) [395] 
458.00ms 544.07ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1 stream 3 tkernel(void) [397] 
458.00ms 544.07ms    (1 1 1)   (1 1 1)   8  0B  0B   -   - TITAN X (Pascal   1 stream 4 tkernel(void) [399] 

Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows. 
SSMem: Static shared memory allocated per CUDA block. 
DSMem: Dynamic shared memory allocated per CUDA block. 
$