TL;DR#

Only with Push/Pop API can you retrieve metrics if auto range is enabled. More specifically:

1
cuptiRangeProfilerPushRange(&pushRangeParams);
2
cuptiRangeProfilerPopRange(&popRangeParams);

Each range is still a kernel. This push and pop only defines the boundry to collect metrics. Any kernels launched outside the region wrapped by push and pop won’t be profiled.

Problem#

According to the sample code in the official document, when the range mode is set to auto range, It is calling pRangeProfilerTarget->PushRange(“name”), which essentially invokes cuptiRangeProfilerPushRange(&pushRangeParams). However, another part of the document indicates that, for the user range, “users can explicitly define ranges using Push/Pop APIs”, which implies that users don’t need to push/pop ranges if they are using auto range. Come on… Which part of the same document should I trust?

Another Potential Way#

I checked another unofficial CUPTI tutorial. The suggested solution is to use Enable/Disable API:

1
// Enable profiling
2
CUpti_Profiler_EnableProfiling_Params enableProfilingParams = {
3
    CUpti_Profiler_EnableProfiling_Params_STRUCT_SIZE};
4
NV_CUPTI_CALL(cuptiProfilerEnableProfiling(&enableProfilingParams));
5

6
// Launch the first kernel (VecAdd)
7
VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
8
cudaDeviceSynchronize();
9

10
// Launch the second kernel (VecSub)
11
VecSub<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_D, N);
12
cudaDeviceSynchronize();
13

14
// Disable profiling
15
CUpti_Profiler_DisableProfiling_Params disableProfilingParams = {
16
    CUpti_Profiler_DisableProfiling_Params_STRUCT_SIZE};
17
NV_CUPTI_CALL(cuptiProfilerDisableProfiling(&disableProfilingParams));

However, the Enable/Disable APIs are Profiling API, which will be deprecated as of CUDA 13.0 and removed in the future. It is recommended to use the Range Profiling API from the header cupti_range_profiler.h.

Experiment#

I speculated 3 possible ways How CUPTI defines the boundry of collection: Push/Pop API, Enable/Disable API or just leave it there(whole program as the boundry). Let’s try it one by one in a simple CUDA program.

1
int main(int argc, char **argv)
2
{
3
    GmpProfiler::getInstance()->init();
4

5
    GmpProfiler::getInstance()->startRangeProfiling();
6

7
    // Can we get the metrics of this kernel?
8
    hello_kernel<<<1,1>>>();
9

10
    GmpProfiler::getInstance()->stopRangeProfiling();
11

12
    cudaDeviceSynchronize();
13
    GmpProfiler::getInstance()->decodeCounterData();
14
    GmpProfiler::getInstance()->printProfilerRanges(outputOption);
15
    GmpProfiler::getInstance()->produceOutput(outputOption);
16

17
    // CUPTI_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL));
18

19
    return 0;
20
}

The GmpProfiler is the customized GPU profiler I built on top of CUPTI. You can think the functions of GmpProfiler as wrappers of CUPTI functions, which can be safely ignored here. After the kernel finishes, GmpProfiler::getInstance()->printProfilerRanges(outputOption); will print the number of ranges in the counter data buffer.

1
collected range size: 0

1
    GmpProfiler::getInstance()->pushRange("hello", GmpProfileType::CONCURRENT_KERNEL);
2
    hello_kernel<<<1,1>>>();
3
    GmpProfiler::getInstance()->popRange("hello", GmpProfileType::CONCURRENT_KERNEL);

1
collected range size: 1

1
    // Enable profiling
2
    CUpti_Profiler_EnableProfiling_Params enableProfilingParams = {CUpti_Profiler_EnableProfiling_Params_STRUCT_SIZE};
3
    cuptiProfilerEnableProfiling(&enableProfilingParams);
4

5
    hello_kernel<<<1,1>>>();
6

7
    // Disable profiling
8
    CUpti_Profiler_DisableProfiling_Params disableProfilingParams = {CUpti_Profiler_DisableProfiling_Params_STRUCT_SIZE};
9
    cuptiProfilerDisableProfiling(&disableProfilingParams);

1
Number of ranges: 0

Conclusion#

Don’t be confused by the official doc, go with

1
cuptiRangeProfilerPushRange(&pushRangeParams);
2
cuptiRangeProfilerPopRange(&popRangeParams);

As we already know, auto range automatically setup each kernel to be a separate range, so the push and pop range here actually means the boundry of data collection. Any kernel launched outside the push and pop calls will not be collected.