Hi! I’m using Jetson Orin NX which has the following specs:
Threads in a warp = 32
Maximum threads allowed in a single SM = 1536
Maximum warps allowed in a single SM = 1536/32 = 48
Total cuda cores in 1 SM = 128 cores = 128/32 warps = 4 warps
So at a time, only 4 warps can be scheduled on the SM while the remaining 44 warps which have been allocated to that SM will have to wait (Scenario 1).
In my case, my kernel was a simple add kernel given here (where N = 100000000):
global void vecAdd(const float* A, const float* B, float* C) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) C[i] = A[i] + B[i];
}
...
...
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C);
According to my understanding let’s say warp 1 has executed the load instruction for the A[i]. Now until the data is being loaded into the allocated register of warp 1, warp 1 will be swapped by another warp (let’s say warp 5). Now warp 5 will be in the execution state while warp 1 will be in the waiting state (scenario 2).
Again according to my understanding the occupancy given by the Nsight-compute is the average number of total active warps during the execution of the kernel. And the warp which has been loaded by the SM, no matter if it’s in the waiting state or execution state, is counted as an active warp.
As shown in the above code, I ran my add example at a block size of 256 threads per block i.e., 8 warps per block and the total number of blocks are 390625. So my total warp count for the add example was 8*390625 = 3125000 warps.
As there are so many warps created during the execution of my program, then still why is my achieved active warps per warps only 40.99 out of 48 (SS of the nsight compute acctached below)
Furthermore I ran another experiment of using 2 cuda streams concurrently. I ran the add kernel twice using 2 different cuda streams hoping that my occupancy will increase and while 2 kernels are running concurrently, the remaining 7 warps slots might get filled by the second stream. I benchmarked it against another code in which I was running the add kernel twice but on a single stream i.e., 1 after the other.
cudaStream_t s1, s2, s3;
cudaStreamCreate(&s1); cudaStreamCreate(&s2); cudaStreamCreate(&s3);
...
...
for (int i = 0; i < benchmark_runs; ++i) {
vecAdd<<<gridSize, blockSize, 0, s1>>>(d_A, d_B, d_C);
vecAdd<<<gridSize, blockSize, 0, s2>>>(d_A, d_B, d_D);
}
...
...
for (int i = 0; i < benchmark_runs; ++i) {
vecAdd<<<gridSize, blockSize, 0, s3>>>(d_A, d_B, d_E);
vecAdd<<<gridSize, blockSize, 0, s3>>>(d_A, d_B, d_F);
}
My single stream code is executing faster than the 2 stream code.
-
Why my achieved occupancy is only 41 warps per SM rather than 48 warps even though I’ve created a lot of warps (according to my understanding a warp in any state idle/waiting/stall/execute is counted as active warp if it’s being grabbed by the SM for execution.) Why the GPU is not utilizing the remaining 7 warp slots?
-
Why was the second stream not able to utilize the remaining available 7 warps and taking more time then a single stream?
-
If the hardware is completely utilized by a single kernel (assuming if it creates more threads than Maximum threads allowed in a single SM * total SMs in GPU, which most of the kernels will usually do), then what is the benifit of streams? As even a basic vector addition will utilize the compelte GPU hardware (assuming vector length N > Maximum threads allowed in a single SM * total SMs in GPU and each thread only computs a single index)