GPU Programming

With i as each thread's index, reading the array every 32 elements as above is slower than reading a[i]. What is the main reason?

float x = a[i * 32];
  1. AAdjacent threads' addresses are 128 bytes apart, so accesses do not coalesce and each thread fetches a separate memory chunk
  2. BBecause the multiplication takes a long time
  3. CBecause reading every 32 elements returns wrong values
Show answer and explanation

Answer: A. Adjacent threads' addresses are 128 bytes apart, so accesses do not coalesce and each thread fetches a separate memory chunk

A float is 4 bytes, so with a[i*32] adjacent threads' addresses are 128 bytes apart. The 32 threads of a warp read 32 separate chunks and use only 4 bytes of each. With a[i], the 32 threads share one 128-byte chunk, so the amount of data moved differs by tens of times. For 2D arrays, the thread index should map to the innermost (contiguous) subscript.

  • A. Correct. Most of each fetched chunk goes unused.
  • B. The multiplication is tiny compared with memory waits.
  • C. The values are correct, just slow.
Practice this topic Back to topic

More GPU Programming questions

  1. When processing particle data with positions x, y and z on a GPU, which layout is better for memory coalescing?
  2. What is a characteristic of CUDA shared memory?
  3. What does CUDA's __syncthreads() do?
  4. In CUDA, __syncthreads() is placed inside "if (threadIdx.x < 16)" so that only some threads reach it. What can happen?
  5. Which describes a bank conflict in CUDA shared memory?