nsys 配置多个进程

Question

我想在 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

但我期待更多这样的事情：Reference

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

（旁注，我在 nvcr 中运行这个例子。io/nvidia/tensorrt:20.12-py3 docker）

Answer 1

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

这是由于低级别 GPU task/context 调度程序（行为）造成的。

过去，调度程序会运行一个 kernel/process/context/task 完成，然后从等待 process/context/task 调度另一个内核。在这种情况下，探查器将描述内核执行而不会重叠。

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

这对于普通内核的一个副作用是，由于时间片，那些似乎同时运行ning 的内核通常需要更长的时间才能运行。对于你的时间延迟内核，时间延迟机制被时间切片“愚弄”（有效地 SM 时钟继续递增）所以它们似乎不再需要运行时间，即使分时分享进行中

（即您已链接的那个）有类似的描述。

nsys 配置多个进程

nsys profile multiple processes

cuda

nvidia

profiler

nsight

mps