NCA Performance Optimization & Debugging 5 — Questions and Answers
Question 1: What is the purpose of the CUDA Occupancy Calculator and when should it be used?
- To determine optimal block size that maximizes active warps per SM given resource usage (Correct answer)
- To estimate PCIe transfer time between host and device
- To predict thermal throttling thresholds for a kernel
- To calculate the number of CUDA cores required for a workload
Correct answer: To determine optimal block size that maximizes active warps per SM given resource usage
The Occupancy Calculator (available via cudaOccupancyMaxPotentialBlockSize API or as a spreadsheet) helps find the thread block size that maximizes SM occupancy given register and shared memory constraints.
Question 2: Which technique allows overlapping CPU computation with GPU kernel execution in CUDA?
- Launching kernels into non-default streams and continuing CPU work before synchronizing (Correct answer)
- Using cudaDeviceSynchronize() immediately after each launch
- Allocating all memory as Unified Memory
- Enabling concurrent kernel execution via cudaSetDeviceFlags
Correct answer: Launching kernels into non-default streams and continuing CPU work before synchronizing
Asynchronous kernel launches into non-default streams return control to the CPU immediately, allowing CPU code to execute in parallel with GPU computation until an explicit sync is needed.
Question 3: A matrix multiplication kernel shows poor performance due to global memory access patterns. Tiling using shared memory improves it. What is the key principle behind this optimization?
- Reusing data loaded into shared memory multiple times reduces total global memory traffic (Correct answer)
- Shared memory has higher bandwidth than L2 cache
- Tiling increases the number of concurrent thread blocks
- Global memory latency is eliminated when tile size matches warp size
Correct answer: Reusing data loaded into shared memory multiple times reduces total global memory traffic
Tiling loads a sub-matrix into shared memory once and reuses it for all dot products in that tile, converting many global reads into fast on-chip shared memory accesses.
Question 4: In Nsight Systems, a developer sees that GPU utilization drops to 0% for several milliseconds between two kernels. Which action best eliminates this gap?
- Pipeline the workload using CUDA streams so the next kernel launches before the previous one completes (Correct answer)
- Increase the number of threads in each kernel
- Use persistent kernels that never return to the host
- Enable CUDA MPS to share the GPU across processes
Correct answer: Pipeline the workload using CUDA streams so the next kernel launches before the previous one completes
Pipelining with CUDA streams lets the CPU enqueue the next kernel while the current one is still running, eliminating idle gaps caused by sequential launch and synchronization patterns.
Question 5: What is the consequence of shared memory bank conflicts on GPU performance?
- Accesses to the same bank by different threads in a warp are serialized, reducing effective bandwidth (Correct answer)
- The SM must flush the shared memory cache between kernel launches
- Threads stall waiting for L2 cache arbitration to resolve
- Register file access becomes a bottleneck instead
Correct answer: Accesses to the same bank by different threads in a warp are serialized, reducing effective bandwidth
When multiple threads in a warp access different addresses within the same shared memory bank, those accesses are serialized into multiple cycles, reducing the effective shared memory throughput.
Question 6: A CUDA program produces different numerical results between CPU and GPU implementations of the same floating-point computation. What is the most likely explanation?
- GPUs use a different order of floating-point operations, and floating-point arithmetic is not associative (Correct answer)
- The GPU rounds all results to single precision regardless of the declared type
- CUDA uses a non-IEEE 754 floating-point format for speed
- The host and device share different endianness
Correct answer: GPUs use a different order of floating-point operations, and floating-point arithmetic is not associative
Floating-point operations are not associative, so different execution orders (parallelism, fused multiply-add, etc.) on GPU vs CPU yield slightly different rounding results—this is expected IEEE 754 behavior.
Question 7: Which flag passed to nvcc enables generation of PTX code for future GPU architectures not yet released at compile time?
- -arch=compute_XY without -code=sm_XY to embed PTX (Correct answer)
- -use_fast_math
- -maxrregcount
- -ptxas-options=-v
Correct answer: -arch=compute_XY without -code=sm_XY to embed PTX
Compiling with -arch=compute_XY (virtual architecture only) embeds PTX bytecode that the CUDA driver can JIT-compile to native ISA for any future GPU compatible with that compute capability.
NCA Performance Optimization & Debugging 5
What is the purpose of the CUDA Occupancy Calculator and when should it be used?