Profile 2 kernels at once

Developer Tools Nsight Compute
1

I have 2 persistent kernels which will run on the SM on the same time. They need to operate together. Now I want to profile both kernels by using ncu. Is there any way for me to profile them in a same launch?

I’ve tried to only profile one kernel, but this operation will prevent the other kernel from launching, I think its weird

2

Hi, @half-0

Do you have a mini-repro for us to try ?
Can you tell which Nsight Compute/driver and the command line you use ?

3

on A100

function collect_ncu() {
  report=$1; shift
  kernel_name=$1; shift

  target_cmd=( "$@" )
  if [[ "${target_cmd[0]}" == "python" || "${target_cmd[0]}" == "python3" ]]; then
    target_cmd[0]="$VIRTUAL_ENV/bin/${target_cmd[0]}"
  fi

  sudo env \
    VIRTUAL_ENV="$VIRTUAL_ENV" \
    PATH="$VIRTUAL_ENV/bin:$PATH" \
    LD_LIBRARY_PATH="$LD_LIBRARY_PATH" \
    /usr/local/cuda-12.8/bin/ncu \
      --set full \
      --force-overwrite \
      --target-processes all \
      --replay-mode application \
      --kernel-name "$kernel_name" \
      --kernel-name-base function \
      --launch-skip 0 \
      --launch-skip-before-match 0 \
      --launch-count 1 \
      --kill no \
      --filter-mode global \
      --warp-sampling-interval auto \
      --warp-sampling-max-passes 5 \
      --warp-sampling-buffer-size 33554432 \
      --profile-from-start 1 \
      --cache-control all \
      --clock-control base \
      --apply-rules yes \
      --import-source no \
      --check-exit-code yes \
      --section MemoryWorkloadAnalysis_Chart \
      --section MemoryWorkloadAnalysis_Tables \
      --metrics group:memory__chart,group:memory__shared_table,group:memory__first_level_cache_table,group:memory__l2_cache_table,group:memory__l2_cache_evict_policy_table,group:memory__dram_table \
      --export "$report" \
      -- "${target_cmd[@]}"
}

collect_ncu "$1" "$2" "${@:3}"

I use this bash to profile a single kernel. However, as I need to launch 2 kernels at the same time:

    kernel_0<<<dim3(54, 1, 1), dim3(128, 1, 1), 98304 /*smem*/, 0>>>(
        runtime_config);

    kernel_1<<<dim3(54, 1, 1), dim3(128, 1, 1), 98304 /*smem*/, 1>>>(
        runtime_config);

only kernel_0 was launched no matter which kernel I profile.

if I swap the launch order, then only kernel_1 will be launched.

As I need 2 kernels to coorporate, the program is stucked

5

Can anyone help me? I still can’t profile it though I’ve used all methods I can use

6

Please define ranges in your code and then profile with range replay. 2. Profiling Guide — NsightCompute 12.9 documentation