消えたCUDA関連の旧ブログ記事を復元するひとり Advent Calendar 2024の記事です。

CUDA Profiler Controlとは

CUDAのRuntime APIやカーネル関数のプロファイルを取りたいときに使うものです．
nvprofのAPI版みたいな気持ちでしょうか
CUDAのすべての呼ばれる関数ではなく，一部の必要な関数のプロファイリングを行いたいときに便利です．
Referenceはあるのですが，configファイルの書き方などがあまりネットに転がっていないように思えたのでサンプルを書いておきます．

Profile Controlサンプル

コード

#include <iostream>
#include <cuda_profiler_api.h>

constexpr std::size_t N = 1lu << 13;
constexpr std::size_t block_size = 1lu << 7;

__global__ void kernel(unsigned long* const a) {
  const auto tid = blockIdx.x + blockDim.x * blockIdx.x;

  a[tid] = tid;
}

int main() {
  cudaProfilerInitialize("profile.conf", "profile.csv", cudaCSV);
  unsigned long *da;
  unsigned long *ha;
  cudaMalloc(&da, sizeof(unsigned long) * N);
  cudaMallocHost(&ha, sizeof(unsigned long) * N);

  // カーネル関数のプロファイル結果だけをCSVで出力
  cudaProfilerStart();
  kernel<<<(N + block_size - 1) / block_size, block_size>>>(da);
  cudaProfilerStop();

  cudaMemcpy(ha, da, sizeof(unsigned long) * N, cudaMemcpyDefault);

  cudaFree(da);
  cudaFreeHost(ha);
}

profile.conf

active_warps
gridsize3d
threadblocksize
dynsmemperblock
stasmemperblock
regperthread
memtransfersize
memtransferdir
memtransferhostmemtype
streamid
cacheconfigrequested
cacheconfigexecuted
countermodeaggregate
enableonstart 0
active_warps
active_cycles

これはPyTorchのコードを参考にしました．
表示できる項目一覧はCOMPUTE COMMAND LINE PROFILER User Guideに書かれています．

結果

こうすることで目的のCSVファイルを得ます． 

# CUDA_PROFILE_LOG_VERSION 2.0
# CUDA_DEVICE 0 GeForce RTX 2080
# CUDA_CONTEXT 1
# CUDA_PROFILE_CSV 1
# TIMESTAMPFACTOR 15b136796fe02291
method,gputime,cputime,gridsizeX,gridsizeY,gridsizeZ,threadblocksizeX,threadblocksizeY,threadblocksizeZ,dynsmemperblock,stasmemperblock,regperthread,occupancy,streamid,cacheconfigexecuted,cacheconfigrequested,active_warps,active_cycles,memtransfersize,memtransferdir,memtransferhostmemtype
_Z6kernelPm,3.296,73.355,64,1,1,128,1,1,0,0,16,1.000,1,0,0,0,0