NCA Performance Optimization & Debugging 3 — Questions and Answers
Question 1: What is the primary purpose of CUDA shared memory in performance optimization?
- To serve as a software-managed cache reducing repeated global memory accesses (Correct answer)
- To increase the number of concurrent thread blocks
- To enable inter-GPU communication
- To store kernel code for faster instruction fetch
Correct answer: To serve as a software-managed cache reducing repeated global memory accesses
Shared memory is on-chip SRAM that threads in a block can use to cache frequently accessed data, dramatically reducing costly global memory transactions.
Question 2: A CUDA kernel uses 64 registers per thread on an SM that supports 65,536 total registers. With 1024 threads per block, how many blocks can run concurrently per SM?
- 1 (Correct answer)
- 2
- 4
- 8
Correct answer: 1
64 registers × 1024 threads = 65,536 registers per block, which exactly exhausts the SM's register budget, allowing only 1 block concurrently.
Question 3: In Nsight Systems, what does a gap between consecutive kernel launches in the GPU timeline most likely indicate?
- CPU-side launch overhead or synchronization stall between kernels (Correct answer)
- The GPU is running out of global memory
- Thermal throttling is reducing GPU clock speed
- Warp divergence is delaying kernel completion
Correct answer: CPU-side launch overhead or synchronization stall between kernels
Gaps in the GPU timeline between kernels reveal that the CPU is not keeping the GPU fed—often due to synchronization calls, data preparation, or excessive launch overhead.
Question 4: Which memory type in CUDA is best suited for read-only data broadcast to all threads in a warp simultaneously?
- Constant memory (Correct answer)
- Global memory
- Local memory
- Texture memory with 2D filtering
Correct answer: Constant memory
Constant memory is cached and optimized for scenarios where all threads in a warp read the same address simultaneously, providing broadcast with very low latency.
Question 5: What happens to CUDA kernel performance when thread blocks are too small (e.g., 32 threads)?
- Poor occupancy limits latency hiding and underutilizes the SM (Correct answer)
- Memory bandwidth increases beyond hardware limits
- Shared memory bank conflicts become more frequent
- Register pressure increases exponentially
Correct answer: Poor occupancy limits latency hiding and underutilizes the SM
Very small blocks reduce the number of warps per block, limiting the SM's ability to hide latency through warp switching and leaving execution units idle.
Question 6: A developer uses __syncthreads() inside a conditional branch where not all threads enter. What is the consequence?
- Undefined behavior or deadlock because all threads in a block must reach the barrier (Correct answer)
- Only threads in the branch synchronize; others continue freely
- The kernel automatically skips the synchronization for inactive threads
- Performance improves because fewer threads participate in the barrier
Correct answer: Undefined behavior or deadlock because all threads in a block must reach the barrier
__syncthreads() requires ALL threads in a block to reach it; placing it inside a divergent branch can cause deadlock or undefined behavior since some threads may never arrive.
Question 7: Which Nsight Compute section reveals whether a kernel is limited by compute throughput versus memory bandwidth?
- Roofline chart in the Speed Of Light section (Correct answer)
- Warp State Statistics
- Source Counters
- Memory Workload Analysis tables
Correct answer: Roofline chart in the Speed Of Light section
The Roofline chart plots achieved FLOPS against arithmetic intensity and visually shows whether the kernel is compute-bound or memory-bandwidth-bound relative to hardware peaks.
NCA Performance Optimization & Debugging 3
What is the primary purpose of CUDA shared memory in performance optimization?