0

我想在 Nvidia GPU 上试验 MPS，因此我希望能够分析两个并行运行的进程。对于现已弃用的 nvprof，曾经有一个选项“--profile-all-processes”。nsys 是否有等价物？

我尝试在MPS OFF的情况下生成多个报告，并使用此代码在同一时间轴上导入它们（来自此问题）：

#include <stdio.h>
#include <stdlib.h>

#define MAX_DELAY 30

#define cudaCheckErrors(msg) \
  do { \
    cudaError_t __err = cudaGetLastError(); \
    if (__err != cudaSuccess) { \
        fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
            msg, cudaGetErrorString(__err), \
            __FILE__, __LINE__); \
        fprintf(stderr, "*** FAILED - ABORTING\n"); \
        exit(1); \
    } \
  } while (0)


#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL

unsigned long long dtime_usec(unsigned long long start){

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}

#define APPRX_CLKS_PER_SEC 1000000000ULL
__global__ void delay_kernel(unsigned seconds){

  unsigned long long dt = clock64();
  while (clock64() < (dt + (seconds*APPRX_CLKS_PER_SEC)));
}

int main(int argc, char *argv[]){
  cudaSetDevice(0);
  unsigned delay_t = 10; // seconds, approximately
  unsigned delay_t_r;
  if (argc > 1) delay_t_r = atoi(argv[1]);
  if ((delay_t_r > 0) && (delay_t_r < MAX_DELAY)) delay_t = delay_t_r;
  unsigned long long difft = dtime_usec(0);
  for (int i = 0; i < 3;i++) {
    delay_kernel<<<1,1>>>(delay_t);
    cudaDeviceSynchronize();
  }
  cudaCheckErrors("kernel fail");
  difft = dtime_usec(difft);
  printf("kernel duration: %fs\n", difft/(float)USECPSEC);
  cudaFree(0);
  return 0;
}

这个脚本：

nvcc -o t1034 t1034.cu

nsys profile -o rep1 -w true -t cuda,nvtx,osrt,cudnn,cublas -s none -f true -x true ./t1034 &
nsys profile -o rep2 -w true -t cuda,nvtx,osrt,cudnn,cublas -s none -f true -x true ./t1034

然后我打开 rep1.qdrep 并将 rep2.qdrep 添加到其中，这会产生以下时间线：Nsys Timeline

但我预计会更像这样：参考

难道我做错了什么 ？这是正确的结果吗？

（旁注，我在 nvcr.io/nvidia/tensorrt:20.12-py3 docker 中运行此示例）

4

1 回答 1

1

我想你的问题是为什么来自不同进程的内核似乎重叠，即使 MPS 已关闭。

其原因是由于低级 GPU 任务/上下文调度程序（行为）。

过去调度程序会运行一个内核/进程/上下文/任务直到完成，然后从某个等待的进程/上下文/任务中调度另一个内核。在这种情况下，分析器将描述没有重叠的内核执行。

最近（假设是在 2015 年之后的某个时间创建您的参考演示文稿时），GPU 调度程序在各种较新的 GPU 和较新的 CUDA 版本上切换到时间片。这意味着从概要分析器的角度来看，在较高级别上，任务似乎“同时”运行，即使 MPS 已关闭。来自进程 1 的内核 A 不一定被允许运行到完成，在上下文调度程序在其轨道中停止该内核、执行上下文切换并允许来自进程 2 的内核 B 开始执行时间片之前。

对于普通内核来说，这样做的一个副作用是，由于时间分片，那些似乎同时运行的内核通常需要更长的时间才能运行。对于您的时间延迟内核，时间延迟机制被时间片“欺骗”（实际上 SM 时钟继续递增），因此即使时间片共享是，它们似乎也不会花费更长的运行时间继续。

这个答案（即你已经链接的那个）有类似的描述。

于 2021-10-15T13:30:25.130 回答