I85PREP
⌘ K
NCA Foundations study materials
NCA Foundations study resources

NCA Foundations practice questions

One question at a time. Build your NCA Foundations understanding.

First 50 of 200 questions. Each one lists the choices, the correct answer, and an explanation. The checker at the top of this page saves progress in this browser when it is available.

Continue in the practice tool above to load more questions, 50 at a time.

Question 1NVIDIA AI Computing Platform and Architecture Fundamentals

Which NVIDIA hardware generation introduced Tensor Cores to accelerate deep learning matrix operations?

  1. Volta Correct answer.
  2. Pascal
  3. Maxwell
  4. Kepler

Explanation

Tensor Cores first appeared in the Volta architecture, specifically in the V100 data center GPU. They perform fused mixed-precision matrix multiply-accumulate operations in hardware, greatly increasing throughput for training and inference workloads compared with earlier designs.

Question 2NVIDIA AI Computing Platform and Architecture Fundamentals

In NVIDIA GPU terminology, what does a Streaming Multiprocessor (SM) primarily contain?

  1. A single CUDA core with private DRAM
  2. A group of CUDA cores plus shared memory and schedulers Correct answer.
  3. Only the memory controller and L2 cache slice
  4. The host CPU interface and PCIe controller

Explanation

An SM is the core building block of an NVIDIA GPU and bundles many CUDA cores together with warp schedulers, registers, and shared memory. Threads are grouped into warps that execute on the SM's resources, making the SM the unit where computation and on-chip data sharing occur.

Question 3NVIDIA AI Computing Platform and Architecture Fundamentals

Which NVIDIA software layer translates CUDA and other parallel programming models so they can run on GPUs from multiple vendors?

  1. CUDA Toolkit
  2. SYCL implementation Correct answer.
  3. OpenCL runtime
  4. NVIDIA HPC SDK

Explanation

SYCL is an open standard for heterogeneous programming, and its implementations can target GPUs from different vendors. CUDA and most NVIDIA toolkits are tied specifically to NVIDIA hardware, so a cross-vendor translation layer is needed to run portable code elsewhere.

Question 4NVIDIA AI Computing Platform and Architecture Fundamentals

What is the main purpose of NVLink in a multi-GPU system?

  1. To convert display output to HDMI
  2. To supply auxiliary power to the GPU board
  3. To provide high-bandwidth direct GPU-to-GPU communication Correct answer.
  4. To connect the GPU to external storage arrays

Explanation

NVLink is a high-speed interconnect that lets GPUs exchange data directly at much higher bandwidth than PCIe. This reduces communication bottlenecks in multi-GPU training and enables faster model-parallel and data-parallel workloads.

Question 5NVIDIA AI Computing Platform and Architecture Fundamentals

Which memory type on an NVIDIA GPU is shared by all threads within a single thread block and offers low latency?

  1. Local memory
  2. Global memory
  3. Constant memory
  4. Shared memory Correct answer.

Explanation

Shared memory is an on-chip scratchpad visible to all threads in a block, allowing fast cooperation and data reuse. Global memory is larger but off-chip and slower, while local and constant memory serve different scoping purposes.

Question 6NVIDIA AI Computing Platform and Architecture Fundamentals

What does the term 'warp' refer to in NVIDIA GPU execution?

  1. A cluster of multiple SMs
  2. A fixed group of 32 threads executed together Correct answer.
  3. A unit of GPU memory allocation
  4. A single thread with the highest priority

Explanation

A warp is the fundamental scheduling unit on NVIDIA GPUs, conventionally containing 32 threads that execute the same instruction in lockstep. Warp-level execution explains concepts such as divergence and coalesced memory access.

Question 7NVIDIA AI Computing Platform and Architecture Fundamentals

Which NVIDIA platform is designed to deliver unified memory addressing between CPU and GPU to simplify programming?

  1. NVIDIA GPUDirect Storage
  2. NVIDIA CUDA Unified Memory Correct answer.
  3. NVIDIA TensorRT
  4. NVIDIA NVSwitch

Explanation

CUDA Unified Memory creates a single address space shared by CPU and GPU, letting the system migrate data automatically. This simplifies code because programmers no longer must manually copy buffers between host and device.

Question 8NVIDIA AI Computing Platform and Architecture Fundamentals

Which NVIDIA GPU feature allows a single physical GPU to be partitioned into multiple isolated instances for separate workloads?

  1. Dynamic boost
  2. GPU Direct
  3. Graphics interop
  4. Multi-Instance GPU (MIG) Correct answer.

Explanation

MIG partitions a supported data center GPU into isolated instances, each with dedicated memory, cache, and compute resources. This isolation improves quality of service when multiple users or jobs share one physical GPU.

Question 9NVIDIA AI Computing Platform and Architecture Fundamentals

A developer wants to accelerate a deep learning model on an NVIDIA GPU using a high-level API that automatically builds and optimizes computation graphs. Which library fits best?

  1. NVIDIA NCCL
  2. NVIDIA DALI
  3. NVIDIA TensorRT Correct answer.
  4. NVIDIA cuBLAS

Explanation

TensorRT is an inference optimization library that builds optimized execution plans for trained networks, applying layer fusion, precision calibration, and kernel selection. cuBLAS provides basic linear algebra, DALI handles data loading, and NCCL handles multi-GPU communication.

Question 10NVIDIA AI Computing Platform and Architecture Fundamentals

Why does coalesced global memory access improve performance on NVIDIA GPUs?

  1. It combines adjacent thread requests into fewer memory transactions Correct answer.
  2. It moves data into registers automatically
  3. It increases the clock frequency of the SM
  4. It disables the L2 cache to reduce overhead

Explanation

When threads in a warp access contiguous, aligned addresses, the memory system can service them with fewer transactions. This raises effective bandwidth because each transaction delivers more useful data per request.

Question 11NVIDIA AI Computing Platform and Architecture Fundamentals

During training, a model's batch size is doubled while the learning rate is kept constant, and throughput per GPU rises but convergence degrades. Which explanation is most consistent with GPU behavior?

  1. Larger batches force the GPU into lower-precision arithmetic automatically
  2. Larger batches reduce shared memory usage, slowing the kernels
  3. Larger batches reduce gradient noise, so the fixed learning rate becomes too small relative to the batch Correct answer.
  4. Larger batches exceed the number of CUDA cores and cause numerical overflow

Explanation

Increasing batch size lowers the variance of the gradient estimate, which often requires a correspondingly larger learning rate to maintain progress. Keeping the rate fixed can therefore slow convergence even though hardware utilization improves.

Question 12NVIDIA AI Computing Platform and Architecture Fundamentals

A kernel shows low occupancy and the profiler reports high register usage per thread. What is the most direct remedy to consider?

  1. Move the kernel to the CPU for execution
  2. Disable error checking in the kernel
  3. Increase the number of blocks launched without changing the kernel
  4. Reduce per-thread register pressure so more warps can reside on the SM Correct answer.

Explanation

Occupancy is limited by the resources each thread consumes, including registers. Lowering register usage allows more warps to be resident on an SM, which can hide latency and improve throughput.

Question 13NVIDIA AI Computing Platform and Architecture Fundamentals

Which NVIDIA interconnect technology enables GPUs to communicate directly with remote GPU memory across nodes in a cluster?

  1. SATA
  2. PCIe Gen3 x16 only
  3. NVLink with GPUDirect RDMA Correct answer.
  4. USB 3.0

Explanation

GPUDirect RDMA lets network adapters move data directly into GPU memory, and combined with NVLink within a node it supports fast cross-node GPU communication. This avoids extra copies through host memory and reduces latency in distributed training.

Question 14NVIDIA AI Computing Platform and Architecture Fundamentals

An inference service must minimize latency for a fixed model on an NVIDIA GPU. Which step is most appropriate?

  1. Disable all GPU memory usage
  2. Train the model from scratch on each request
  3. Use TensorRT to build an optimized engine with layer fusion and precision calibration Correct answer.
  4. Increase the batch size to the maximum regardless of latency targets

Explanation

TensorRT optimizes trained networks for inference by fusing layers, selecting efficient kernels, and calibrating reduced precision where acceptable. These steps lower per-request latency without retraining the model.

Question 15NVIDIA AI Computing Platform and Architecture Fundamentals

What advantage does half precision (FP16) offer over FP32 on Tensor Core-equipped GPUs?

  1. Higher throughput for matrix operations with reduced memory footprint Correct answer.
  2. Elimination of the need for any memory transfers
  3. Support for larger integer ranges
  4. Guaranteed higher numerical accuracy

Explanation

FP16 halves the storage per value and lets Tensor Cores execute matrix multiply-accumulate operations at higher rates than FP32. The tradeoff is reduced precision, which often requires techniques such as loss scaling during training.

Question 16NVIDIA AI Computing Platform and Architecture Fundamentals

A team runs a multi-GPU training job and observes that communication dominates runtime. Which library is designed to optimize collective operations such as all-reduce across GPUs?

  1. NCCL Correct answer.
  2. NPP
  3. cuFFT
  4. cuSPARSE

Explanation

NCCL provides optimized implementations of collective communication primitives like all-reduce, broadcast, and all-gather across GPUs and nodes. Using NCCL typically reduces synchronization overhead compared with naive communication schemes.

Question 17NVIDIA AI Computing Platform and Architecture Fundamentals

Which statement best describes the role of the L2 cache in an NVIDIA GPU?

  1. It stores the operating system kernel
  2. It is private to each CUDA core and stores only constants
  3. It replaces global memory entirely
  4. It is shared across SMs and buffers data between the SMs and global memory Correct answer.

Explanation

The L2 cache is a large on-chip memory shared by all SMs, sitting between the compute units and global DRAM. It reduces average memory latency and can improve bandwidth for reused data.

Question 18NVIDIA AI Computing Platform and Architecture Fundamentals

A developer needs to run many small independent jobs on one data center GPU with strict memory isolation between them. Which capability is most suitable?

  1. GPU boost clocks
  2. MIG instances Correct answer.
  3. A single large CUDA context shared by all jobs
  4. CUDA streams within one process

Explanation

MIG creates hardware-isolated GPU instances, each with its own memory slice and compute resources, so one job cannot consume another's memory. Shared contexts or streams do not provide that isolation.

Question 19NVIDIA AI Computing Platform and Architecture Fundamentals

Which change is most likely to improve performance when a kernel is limited by global memory bandwidth rather than compute?

  1. Adding more arithmetic instructions per thread
  2. Increasing data reuse through shared memory tiling Correct answer.
  3. Raising the SM clock frequency only
  4. Increasing the number of kernel arguments

Explanation

When bandwidth-bound, the goal is to move fewer bytes from global memory, and tiling data into shared memory lets threads reuse values. Adding arithmetic or clock changes does not reduce the memory traffic that limits the kernel.

Question 20NVIDIA AI Computing Platform and Architecture Fundamentals

What is the primary function of the NVIDIA CUDA Toolkit?

  1. Providing a relational database engine
  2. Providing compilers, libraries, and tools for building GPU-accelerated applications Correct answer.
  3. Serving as a container runtime for Kubernetes
  4. Managing physical data center power distribution

Explanation

The CUDA Toolkit bundles the NVCC compiler, runtime libraries, debugging and profiling tools, and math libraries used to develop GPU-accelerated software. It is a development environment rather than infrastructure or database software.

Question 21NVIDIA AI Computing Platform and Architecture Fundamentals

In a CUDA program, what distinguishes host code from device code?

  1. Both run exclusively on the GPU
  2. Host code runs on the CPU and device code runs on the GPU Correct answer.
  3. Host code runs on the GPU and device code runs on the CPU
  4. Both run exclusively on the CPU

Explanation

CUDA uses the host, typically the CPU, to orchestrate execution and the device, the GPU, to run parallel kernels. Data is transferred between host and device memory as needed, and kernels are launched from host code.

Question 22NVIDIA AI Computing Platform and Architecture Fundamentals

Which factor most directly limits the maximum number of resident warps on a single SM?

  1. The speed of the host CPU
  2. The number of PCIe lanes in the system
  3. The size of the display framebuffer
  4. Available registers, shared memory, and warp slots per SM Correct answer.

Explanation

Occupancy is bounded by hardware limits such as the maximum warps per SM and the per-thread register and per-block shared memory budgets. Host CPU speed, display memory, and PCIe lane counts do not determine how many warps fit on an SM.

Question 23NVIDIA AI Computing Platform and Architecture Fundamentals

A profiling run shows warp divergence in a branch-heavy kernel. Which strategy best reduces its performance impact?

  1. Reducing the number of blocks in the grid
  2. Moving the branch into a separate kernel launch for every thread
  3. Increasing the number of threads per block without changing the code
  4. Restructuring the code so threads within a warp follow the same branch path Correct answer.

Explanation

Divergence occurs when threads in a warp take different paths, forcing serialized execution of each branch. Reorganizing data or logic so that threads in a warp share the same path restores parallel efficiency.

Question 24NVIDIA AI Computing Platform and Architecture Fundamentals

Why is NVLink generally preferred over PCIe for tightly coupled multi-GPU model parallelism?

  1. It eliminates the need for any synchronization
  2. It reduces the GPU power draw significantly
  3. It offers higher peer-to-peer bandwidth and lower latency between GPUs Correct answer.
  4. It provides more storage capacity per GPU

Explanation

Model parallelism requires frequent exchange of activations or gradients between GPUs, so interconnect bandwidth and latency dominate. NVLink provides substantially higher GPU-to-GPU bandwidth than PCIe, reducing communication stalls.

Question 25NVIDIA AI Computing Platform and Architecture Fundamentals

A team observes that a mixed-precision training run diverges with FP16 but succeeds in FP32. Which remedy is most appropriate?

  1. Increase the batch size to compensate
  2. Apply loss scaling to preserve small gradient magnitudes Correct answer.
  3. Switch the optimizer to SGD without changing anything else
  4. Disable Tensor Cores permanently

Explanation

FP16 has a limited exponent range, so small gradients can underflow to zero. Loss scaling multiplies the loss before backpropagation and unscales gradients afterward, preserving values that would otherwise be lost.

Question 26NVIDIA AI Computing Platform and Architecture Fundamentals

Which NVIDIA hardware unit is responsible for executing the arithmetic and logic instructions of a CUDA thread?

  1. Copy engine
  2. Texture unit
  3. Display controller
  4. CUDA core Correct answer.

Explanation

CUDA cores are the scalar arithmetic pipelines inside a streaming multiprocessor that perform integer and floating-point operations for the threads assigned to them. Texture units sample and filter textures, and copy engines move data between host and device memory rather than executing thread instructions.

Question 27NVIDIA AI Computing Platform and Architecture Fundamentals

What is the primary purpose of the NVCC compiler in the CUDA programming model?

  1. It compiles CUDA source into host and device code Correct answer.
  2. It schedules kernels across multiple GPUs
  3. It profiles kernel memory bandwidth
  4. It manages GPU memory allocation at runtime

Explanation

NVCC separates CUDA source into host code and device code, compiling device code for the target GPU architecture and forwarding host code to a standard host compiler. The resulting fat binary is then linked into the application.

Question 28NVIDIA AI Computing Platform and Architecture Fundamentals

A developer launches a kernel with a grid of 100 blocks and 256 threads per block. On a GPU whose threads are grouped into warps of 32, how many warps does the entire grid contain?

  1. 3200
  2. 25600
  3. 800 Correct answer.
  4. 100

Explanation

Each block of 256 threads forms 256 divided by 32, which is 8 warps. Multiplying by 100 blocks gives 800 warps across the whole grid, since warp count scales with total threads divided by the warp size.

Question 29NVIDIA AI Computing Platform and Architecture Fundamentals

Which CUDA memory space is shared by all threads within a single thread block and is typically used to hold data reused across those threads?

  1. Global memory
  2. Local memory
  3. Shared memory Correct answer.
  4. Constant memory

Explanation

Shared memory is an on-chip scratchpad scoped to a thread block, so threads in the same block can cooperate through it at low latency. Global memory is visible to all threads but resides off-chip, while constant and local memory serve different access patterns.

Question 30NVIDIA AI Computing Platform and Architecture Fundamentals

A kernel performs a reduction across a large array and currently has each thread read one element from global memory, then uses shared memory for the tree reduction. Profiling shows global memory transactions dominate. Which change most directly reduces global memory traffic in the reduction stage?

  1. Add more threads per block so each thread handles exactly one element
  2. Increase the number of blocks so more elements are processed in parallel
  3. Store the partial results in constant memory instead of shared memory
  4. Have each thread load several elements and accumulate them before the shared-memory tree Correct answer.

Explanation

Having each thread accumulate multiple elements reduces the number of threads and thus the number of global loads and shared-memory tree levels needed. This serial multi-element accumulation lowers memory transactions per output while keeping the reduction correct.

Question 31GPU Architecture and Parallel Computing Concepts

In the CUDA execution model, what is the primary role of a warp?

  1. A group of 32 threads that a streaming multiprocessor schedules and executes together Correct answer.
  2. A cluster of streaming multiprocessors that share one L1 cache
  3. A software library that manages host-device memory transfers
  4. A group of threads that share a single private register file

Explanation

A warp is the fundamental scheduling unit on NVIDIA GPUs, consisting of 32 threads that the warp scheduler issues together. Threads within a warp execute the same instruction stream, which is why divergence within a warp reduces efficiency.

Question 32GPU Architecture and Parallel Computing Concepts

Which memory space on an NVIDIA GPU is shared by all threads in a block but is private to that block?

  1. Shared memory Correct answer.
  2. Constant memory
  3. Local memory
  4. Global memory

Explanation

Shared memory is allocated per thread block and is visible to every thread within that block, enabling low-latency data exchange. It is not shared across different blocks, and its lifetime matches the block's execution.

Question 33GPU Architecture and Parallel Computing Concepts

A kernel is launched with a grid of 8 blocks and each block has 256 threads. How many threads execute in total?

  1. 2048 Correct answer.
  2. 4096
  3. 1024
  4. 256

Explanation

The total thread count equals the number of blocks multiplied by the threads per block. Here 8 blocks times 256 threads per block gives 2048 threads.

Question 34GPU Architecture and Parallel Computing Concepts

What does the term SIMT describe in GPU architecture?

  1. A single instruction stream applied to a single data element per cycle
  2. Multiple independent instruction streams each with its own program counter per thread
  3. A single instruction issued to multiple threads that each keep their own registers and data Correct answer.
  4. A memory model where all threads read and write one shared address simultaneously

Explanation

SIMT, single instruction multiple thread, means one instruction is issued for a group of threads while each thread maintains its own registers and data. It resembles SIMD but exposes per-thread state and allows independent branching paths.

Question 35GPU Architecture and Parallel Computing Concepts

Which statement best describes the purpose of a streaming multiprocessor (SM) in an NVIDIA GPU?

  1. It manages the PCIe link between the CPU and the GPU
  2. It hosts warps, registers, shared memory, and execution cores that run thread blocks Correct answer.
  3. It stores the compiled kernel binary for the entire device
  4. It converts floating-point values into fixed-point values for the host

Explanation

An SM is the core execution unit that contains warp schedulers, registers, shared memory, and many arithmetic cores. Thread blocks are assigned to SMs, and the SM runs the warps of those blocks concurrently.

Question 36GPU Architecture and Parallel Computing Concepts

Which of the following is a characteristic of global memory on a GPU?

  1. It is a read-only region cached in the texture pipeline
  2. It is a small register file used for per-thread temporaries
  3. It is accessible only by threads in the same block and is cleared at kernel exit
  4. It is accessible by all threads and persists across kernel launches until explicitly freed Correct answer.

Explanation

Global memory is device-wide, visible to every thread in every block, and its contents persist across kernel launches until the allocation is freed. It typically has the largest capacity but higher latency than shared memory.

Question 37GPU Architecture and Parallel Computing Concepts

What is memory coalescing in GPU global memory access?

  1. Compressing data before writing it to global memory
  2. Copying global memory into shared memory before every load
  3. Encrypting memory transactions between the GPU and host
  4. Combining accesses from threads in a warp into fewer memory transactions when addresses are contiguous Correct answer.

Explanation

When threads in a warp access contiguous, aligned addresses, the hardware can service them with fewer memory transactions, improving effective bandwidth. Scattered addresses require more transactions and reduce performance.

Question 38GPU Architecture and Parallel Computing Concepts

Which resource limit directly determines how many warps of a block can reside simultaneously on an SM?

  1. The register and shared memory usage per thread and per block Correct answer.
  2. The clock rate of the display output
  3. The size of the kernel source file
  4. The number of host threads in the process

Explanation

Occupancy is bounded by resources such as registers per thread, shared memory per block, and hardware limits on warps and blocks per SM. Higher per-thread register or per-block shared memory usage reduces the number of resident warps.

Question 39GPU Architecture and Parallel Computing Concepts

Threads in a warp take different branches of an if-else statement. What is the typical hardware behavior?

  1. The compiler removes the conditional so both paths run for all threads
  2. The warp is split into two permanent warps that never recombine
  3. The kernel aborts because divergence is illegal
  4. Both paths execute serially with the inactive threads masked, then the warp reconverges Correct answer.

Explanation

Warp divergence causes the divergent paths to be serialized, with threads not on the active path masked off. After the paths complete, execution reconverges, so divergence costs extra cycles rather than causing an error.

Question 40GPU Architecture and Parallel Computing Concepts

A kernel reads a 2D array with column-major stride across threads in a warp. What performance issue is most likely?

  1. Loss of shared memory capacity across the whole SM
  2. Automatic conversion of global memory to constant memory
  3. Uncoalesced global memory accesses generating many transactions Correct answer.
  4. Increased register spilling to local memory

Explanation

When consecutive threads access addresses separated by a large stride, the warp's accesses are not contiguous, so the hardware issues many separate transactions. This uncoalesced pattern reduces effective bandwidth.

Question 41GPU Architecture and Parallel Computing Concepts

Why is shared memory often used as a staging buffer in tiled matrix multiplication?

  1. It automatically parallelizes the inner loop across blocks
  2. It removes the need for synchronization within a block
  3. It allows each thread to write directly to host memory
  4. It provides low-latency reuse of data loaded once from global memory by many threads Correct answer.

Explanation

Tiling loads a block of data from global memory into shared memory once, then many threads reuse it, cutting global traffic. Shared memory latency is much lower than global, but barriers are still needed to coordinate loads and use.

Question 42GPU Architecture and Parallel Computing Concepts

What is the main effect of increasing occupancy on an SM?

  1. It gives the scheduler more warps to hide memory and execution latency Correct answer.
  2. It reduces the total amount of global memory available
  3. It eliminates warp divergence in all kernels
  4. It guarantees every warp issues an instruction every cycle

Explanation

More resident warps give the warp scheduler additional candidates to issue when others stall, which helps hide latency. Occupancy alone does not guarantee full issue or remove divergence.

Question 43GPU Architecture and Parallel Computing Concepts

Which synchronization primitive ensures all threads in a block reach a point before any proceed?

  1. __ldg
  2. __syncthreads() Correct answer.
  3. atomicAdd
  4. cudaMemcpy

Explanation

__syncthreads() is a block-level barrier that blocks each thread until all threads in the block arrive. It is commonly used after loading shared memory to ensure the data is ready for all threads.

Question 44GPU Architecture and Parallel Computing Concepts

A kernel uses 64 registers per thread and the SM has 65536 registers. Ignoring other limits, how many threads can be resident from register usage alone?

  1. 512
  2. 4096
  3. 2048
  4. 1024 Correct answer.

Explanation

Dividing 65536 registers by 64 registers per thread gives 1024 threads. This register-based limit caps occupancy regardless of other resources until it is exceeded by other constraints.

Question 45GPU Architecture and Parallel Computing Concepts

What is the purpose of a grid-stride loop in a CUDA kernel?

  1. To convert global memory accesses into shared memory accesses
  2. To guarantee that no warp divergence occurs
  3. To force all blocks to execute in lockstep
  4. To let a fixed grid process arrays larger than the total number of launched threads Correct answer.

Explanation

A grid-stride loop advances the index by the total number of threads in the grid each iteration, so a fixed grid can cover arbitrarily large data. It also supports launching a grid sized to the device rather than the data.

Question 46GPU Architecture and Parallel Computing Concepts

Which memory access pattern best enables broadcast within a warp reading a single shared memory address?

  1. Each thread reads a different bank in the same instruction
  2. All threads read the same shared memory address in the same instruction Correct answer.
  3. Threads alternate between shared and global memory reads
  4. Each thread reads consecutive addresses spanning all banks

Explanation

When all threads in a warp read the same shared memory address, the hardware can broadcast the single value to all threads without bank conflicts. Distinct addresses in the same bank would serialize instead.

Question 47GPU Architecture and Parallel Computing Concepts

Why can excessive use of __syncthreads() inside a loop hurt performance?

  1. It increases the register count of every thread
  2. It forces warps to wait, reducing overlap and exposing latency Correct answer.
  3. It disables the L1 cache for the entire kernel
  4. It converts shared memory into local memory

Explanation

Each barrier requires all threads in the block to arrive before any continue, which limits how much independent work warps can overlap. Frequent barriers therefore expose latency that would otherwise be hidden.

Question 48GPU Architecture and Parallel Computing Concepts

What distinguishes constant memory from global memory in typical GPU usage?

  1. Constant memory is read-only from the kernel and benefits from a dedicated cache when all threads read the same address Correct answer.
  2. Constant memory has larger capacity than global memory
  3. Constant memory is writable by all threads and uncached
  4. Constant memory is allocated per block and cleared at kernel exit

Explanation

Constant memory is read-only during kernel execution and is served by a dedicated constant cache that performs well when a warp reads the same address. Global memory is writable and much larger in capacity.

Question 49GPU Architecture and Parallel Computing Concepts

A warp repeatedly issues loads that hit in L1 cache. How does this affect the need to switch warps?

  1. Cache hits increase the latency of each load
  2. Cache hits remove the need for multiple resident warps entirely
  3. Short-latency hits reduce stalls, so fewer switches are needed to keep the pipeline busy Correct answer.
  4. Cache hits force the scheduler to serialize all warps

Explanation

Latency hiding is needed because memory operations stall warps; cache hits shorten that latency, so the scheduler needs fewer alternative warps to keep issuing. The benefit depends on hit rate and the rest of the workload.

Question 50GPU Architecture and Parallel Computing Concepts

Which factor makes a thread block a natural unit for cooperation?

  1. All threads in a block always execute in a single warp
  2. All threads in a block are guaranteed to run on the same SM and can use block barriers Correct answer.
  3. Blocks are guaranteed to run in a fixed order across the device
  4. Blocks share a single register file across the whole GPU

Explanation

A block is scheduled onto one SM, so its threads can share memory and synchronize with block barriers. Different blocks may run anywhere and in any order, so they cannot rely on block-level coordination.