GPU Programming

When processing particle data with positions x, y and z on a GPU, which layout is better for memory coalescing?

  1. AAn array of structures with x, y and z per particle
  2. BArrange all values in random order
  3. CSeparate arrays for x, y and z (a structure of arrays rather than an array of structures)
Show answer and explanation

Answer: C. Separate arrays for x, y and z (a structure of arrays rather than an array of structures)

In an array of structures (AoS), the layout is x, y, z of particle 0, then x, y, z of particle 1 and so on, so when each thread reads x the addresses are 12 bytes apart. In a structure of arrays (SoA), all x values are contiguous, so adjacent threads read adjacent x values and the accesses coalesce. SoA is often better for CPU SIMD for the same reason.

  • A. Adjacent threads' x values end up 12 bytes apart.
  • B. That would prevent coalescing.
  • C. Correct. Threads reading the same component access consecutive addresses.
Practice this topic Back to topic

More GPU Programming questions

  1. What is a characteristic of CUDA shared memory?
  2. What does CUDA's __syncthreads() do?
  3. In CUDA, __syncthreads() is placed inside "if (threadIdx.x < 16)" so that only some threads reach it. What can happen?
  4. Which describes a bank conflict in CUDA shared memory?
  5. A 32 × 32 float tile in shared memory causes bank conflicts when a warp's 32 threads read down a column. What is the common fix?