Как я могу получить доступ к числовым идентификаторам потоков, которые можно увидеть в nvprof, используя cudaStream_t?

Question

В nvprof я могу видеть идентификаторы потока для каждого потока выполнения cuda, который я использую (0, 13, 15 и т. Д.)

Учитывая переменную потока, я хотел бы иметь возможность распечатать идентификатор потока. В настоящее время я не могу найти какой-либо API для этого и приведение cudaStream_t к int или uint не дает разумного идентификатора. sizeof() говорится cudaStream_t 8 байт.

Answer 1

Вкратце: я не знаю способа прямого доступа к этим идентификаторам, но вы можете дать потокам явные имена для профилирования.

---

cudaStream_t является непрозрачным типом "дескриптора ресурса". Дескриптор ресурса - это что-то вроде указателя; поэтому понятно, что идентификатор потока содержится не в самом указателе (дескрипторе), а как-то в том, к чему он относится.

Так как он непрозрачный (без определения того, на что он указывает, предоставленный CUDA) и, как вы указываете, для этого нет прямого API, я не думаю, что вы найдете способ извлечь идентификатор потока из cudaStream_t во время выполнения.

Для этих утверждений, что cudaStream_t дескриптор ресурса и что он непрозрачный, обратитесь к заголовочному файлу CUDA driver_types.h

Однако API расширения NV Tools дает вам возможность "назвать" конкретный поток (или другие ресурсы). Это позволит вам связать определенный поток в исходном коде с определенным именем в профилировщике.

Вот тривиальный рабочий пример:

$ 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.
$