It is a sampled measurement over a time period. For a given time period, it reports what percentage of time one or more GPU kernel(s) was active (i.e. running).
It doesn’t tell you anything about how many SMs were used, or how “busy” the code was, or what it was doing exactly, or in what way it may have been using memory.
The above claim(s) can be verified without too much difficulty using a microbenchmarking-type exercise (see below).
Based on the Nvidia docs, The sample period may be between 1 second and 1/6 second depending on the product. However, the period shouldn’t make much difference on how you interpret the result.
Also, the word “Volatile” does not pertain to this data item in nvidia-smi. You are misreading the output format.
Here’s a trivial code that supports my claim:
#include <stdio.h>
#include <unistd.h>
#include <stdlib.h>
const long long tdelay=1000000LL;
const int loops = 10000;
const int hdelay = 1;
__global__ void dkern(){
long long start = clock64();
while(clock64() < start+tdelay);
}
int main(int argc, char *argv[]){
int my_delay = hdelay;
if (argc > 1) my_delay = atoi(argv[1]);
for (int i = 0; i<loops; i++){
dkern<<<1,1>>>();
usleep(my_delay);}
return 0;
}
On my system, when I run the above code with a command line parameter of 100, nvidia-smi will report 99% utilization. When I run with a command line parameter of 1000, nvidia-smi will report ~83% utilization. When I run it with a command line parameter of 10000, nvidia-smi will report ~9% utilization.
Although this answer is focused on GPU kernels, I have lately noticed that nvidia-smi will also report non-zero GPU utilization when for example cudaMemcpy operations are running (and nothing else). So the above description should be considered a description of reporting with respect to CUDA kernel activity.