CODEWITHSUNDEEP
cudaprofilingnvvpnvprof
travelingbones
travelingbones

Reputation: 8448

How to observe CUDA events and metrics for a subsection of an executable (e.g. only during a kernel execution time)?

I'm familiar with using nvprof to access the events and metrics of a benchmark, e.g., 

nvprof --system-profiling on --print-gpu-trace -o (file name) --events inst_issued1 ./benchmarkname

The 

system-profiling on --print-gpu-trace -o (filename)

command gives timestamps for start time, kernel end times, power, temp and saves the info an nvvp files so we can view it in the visual profiler. This allows us to see what's happening in any section of a code, in particular when a specific kernel is running. My question is this--

Is there a way to isolate the events counted for only a section of the benchmark run, for example during a kernel execution? In the command above, 

--events inst_issued1

just gives the instructions tallied for the whole executable. Thanks!

Upvotes: 9

Views: 9051

Answers (2)

Robert Crovella
Robert Crovella

Reputation: 152164

You may want to read the profiler documentation.

You can turn profiling on and off within an executable. The cuda runtime API for this is:

cudaProfilerStart() 
cudaProfilerStop()

So, if you wanted to collect profile information only for a specific kernel, you could do:

#include <cuda_profiler_api.h>
...

cudaProfilerStart();
myKernel<<<...>>>(...);
cudaProfilerStop();

(instead of a kernel call, the above could be a function or code that calls kernels) Excerpting from the documentation:

When using the start and stop functions, you also need to instruct the profiling tool to disable profiling at the start of the application. For nvprof you do this with the --profile-from-start off flag. For the Visual Profiler you use the Start execution with profiling enabled checkbox in the Settings View.

Also from the documentation for nvprof specifically, you can limit event/metric tabulation to a single kernel with a command line switch:

 --kernels <kernel name>

The documentation gives additional usage possibilities.

The same methodologies are possible with nsight systems and nsight compute.

The CUDA profiler start/stop functions work exactly the same way. The nsight systems documentation explains how to run the profiler with capture control managed by the profiler api:

nsys [global-options] start -c cudaProfilerApi

or for nsight compute:

ncu [options] --profile-from-start off

Likewise, nsight compute can be conditioned via the command line to only profile specific kernels. The primary switch for this is -k to select kernels by name. In repetetive situations, the -c switch can be used to determine the number of the named kernel launches to profile, and the -s switch can be used to skip a number of launches before profiling.

These methodologies don't apply just to events and metrics, but to all profiling activity performed by the respective profilers.

The CUDA profiler API can be used in any executable, and does not require compilation with nvcc.

Upvotes: 20

travelingbones
travelingbones

Reputation: 8448

After looking into this a bit more, it turns out that kernel level information is also given for all kernels (w/o using --kernels and specifying them specifically) by using 

nvprof --events <event names> --metrics <metric names> ./<cuda benchmark>

In fact, it gives output of the form

"Device","Kernel","Invocations","Event Name","Min","Max","Avg"

If a kernel is called multiple times in the benchmark, this allows you to see the Min, Max, Avg of the desired events for those kerne runs. Evidently the --kernels option on Cuda 7.5 Profiler allows each run of each kernel to be specified.

Upvotes: 0

Related Questions