Self-Assessment Quiz: GPU Computing

Twenty questions to confirm the concepts landed — the host/device model, the two roads (OpenACC and CUDA Fortran), and above all the data-movement discipline that decides whether an offload is worth doing. Aim for 16 or more. Reason the code questions by hand; the answer key and a topic map are at the end.

Question 1

A GPU differs from a CPU chiefly in that it is optimized for: - A. Low latency on a single thread of execution - B. High throughput across thousands of simple cores doing the same operation - C. Large caches and branch prediction for irregular code - D. Running the operating system

Question 2

In accelerator programming, the host and the device are: - A. Two names for the same processor - B. The CPU (+ its RAM) and the GPU (+ its own memory), with separate memory spaces - C. Two GPUs in the same machine - D. The compiler and the linker

Question 3

The single most important consequence of host and device having separate memory is: - A. The GPU is always faster - B. Data must be explicitly copied between them before and after a kernel runs - C. You cannot use Fortran on a GPU - D. The CPU is idle during a kernel

Question 4

Which OpenACC directive offloads the following loop to the GPU, running its iterations in parallel? - A. !$omp parallel do - B. !$acc parallel loop - C. !$acc data - D. !$cuf kernel

Question 5

The OpenACC data clause for an array the kernel only reads as input is: - A. copyout - B. create - C. copyin - D. present

Question 6

The purpose of an !$acc data region around a time loop is to: - A. Make the loop run in parallel - B. Keep arrays resident on the device so they are transferred once, not every step - C. Allocate host memory - D. Synchronize CPU threads

Question 7

True or false, with justification: OpenACC directives are structured comments, so a source file with !$acc directives still compiles and runs correctly (on the CPU) with a compiler that ignores them.

Question 8

In CUDA Fortran, a kernel is a subroutine marked: - A. attributes(host) - B. attributes(device) - C. attributes(global) - D. pure

Question 9

A CUDA Fortran kernel must be defined: - A. Inside the main program's contains - B. As a module procedure - C. In a separate file only - D. As an internal function

Question 10

In CUDA Fortran, the correct global index for a thread is: - A. i = blockIdx%x*blockDim%x + threadIdx%x - B. i = (blockIdx%x - 1)*blockDim%x + threadIdx%x - C. i = blockIdx%x + threadIdx%x - D. i = threadIdx%x*gridDim%x + blockIdx%x

Question 11

Why does CUDA Fortran's index formula subtract 1 from blockIdx%x, unlike CUDA C's? - A. Fortran arrays start at 0 - B. CUDA Fortran's threadIdx%x and blockIdx%x are one-based, matching Fortran arrays - C. It is a bug that everyone copies - D. To skip the first block

Question 12

In the launch call k<<<nblocks, tpb>>>(...), the two numbers are: - A. Grid dimensions in x and y - B. Number of blocks, and threads per block - C. Input and output sizes - D. The compute capability

Question 13

Why does a kernel launched with more threads than data elements need an if (i <= n) guard? - A. To make it run faster - B. So the extra threads (whose index exceeds n) do not read or write out of bounds - C. To synchronize the threads - D. To free device memory

Question 14

In a GPU program, the usual performance bottleneck is: - A. The on-device arithmetic - B. The host–device data transfer across the bus - C. The compiler - D. Printing the results

Question 15

A stencil offload transfers the whole field to the device and back every time step. The likely result is: - A. A large speedup - B. The GPU version is slower than the CPU version, because transfers dominate - C. A compile error - D. Exactly the CPU speed

Question 16

True or false, with justification: "My kernel runs 20× faster than the CPU loop" guarantees the whole program runs about 20× faster.

Question 17

Which computation is the best fit for a GPU? - A. A 15-iteration loop where each step depends on the last - B. A recursive traversal with data-dependent branching - C. The same five-point stencil applied to ten million grid cells, swept thousands of times - D. Reading a configuration file

Question 18

"Warp divergence" hurts GPU performance because: - A. The GPU runs out of memory - B. Threads in a lock-step group take different branches, so the hardware runs the paths serially - C. The bus is too slow - D. Fortran is column-major

Question 19

Compared with OpenACC, CUDA Fortran is: - A. More portable, less control - B. NVIDIA-only, more explicit control over kernels and device memory - C. Unable to use device memory - D. A directive-based standard

Question 20

You keep a field resident on the device but must write a frame to disk every 500 steps. The directive that refreshes the host copy without ending the data region is: - A. !$acc end data - B. !$acc update self(field) - C. !$acc parallel loop - D. !$acc create(field)

Answer Key

Q Ans Why
1 B A GPU is a throughput machine — thousands of simple cores doing the same operation on much data.
2 B Host = CPU + RAM; device = GPU + its own memory; the two are separate address spaces.
3 B Separate memory means data must be explicitly copied across the bus before and after a kernel.
4 B !$acc parallel loop offloads the following loop and runs its iterations in parallel on the device.
5 C copyin moves a read-only input to the device with no copy-back.
6 B A data region keeps arrays resident, so they cross the bridge once (in/out) rather than every step.
7 True Directives are comments; a non-OpenACC compiler ignores them and the code runs correctly on the CPU.
8 C attributes(global) marks a kernel — called from the host, run on the device.
9 B Kernels must be module procedures, not internal contains procedures of the program.
10 B CUDA Fortran is one-based, so i = (blockIdx%x - 1)*blockDim%x + threadIdx%x.
11 B threadIdx%x/blockIdx%x are one-based in CUDA Fortran (matching Fortran arrays), so subtract 1.
12 B The launch config is <<<number of blocks, threads per block>>>.
13 B Extra threads have i > n; without the guard they access out of bounds.
14 B The host–device bus is far slower than device memory; transfers usually dominate.
15 B Per-step transfers swamp a cheap stencil, so the GPU version runs slower than the CPU one.
16 False Offloading speeds only the offloaded part; serial work and transfers cap the whole-program speedup (Amdahl).
17 C Large, regular, data-parallel, heavily reused on the device — the ideal GPU workload.
18 B Divergent branches within a warp run serially, so a branchy warp runs at a fraction of peak.
19 B CUDA Fortran is NVIDIA-only and explicit; OpenACC is portable and directive-based.
20 B !$acc update self (a.k.a. update host) refreshes the host copy without ending residency.

Topics to review by question

  • Q1–3 → §35.1 (the GPU, host/device, separate memory).
  • Q4–7 → §35.2 (OpenACC directives, data clauses, data regions).
  • Q8–13 → §35.3 (CUDA Fortran kernels, one-based indexing, launch config, the tail guard).
  • Q14–16 → §35.4 (host–device transfer as the bottleneck; Amdahl still rules).
  • Q17–18 → §35.5 (when GPUs help vs hurt; warp divergence).
  • Q19 → §35.2 vs §35.3 (the OpenACC/CUDA Fortran trade).
  • Q20 → §35.4 (update self for output from a resident field).

Scored below 16? The two ideas most worth rereading are the host/device split with its one commandment — move data across the bridge as rarely as possible (Q3, 6, 14, 15, 20) — and the one-based CUDA Fortran thread index (Q10, 11). They are where correctness and performance actually live in this chapter.