4 The Map pattern

In this chapter, the map pattern is presented as the simplest and most universal parallel building block: the same operation is applied independently to many data elements, so each thread can work on one output without coordinating with any other thread. The core idea is that a map has no cross-output dependence, and in a pure map each output reads only the input at the same index. That independence is what makes map “embarrassingly parallel” and ideal for GPUs, CPUs with SIMD or OpenMP, and even distributed systems such as MapReduce.

The chapter then shows how to recognize map in practice using a checklist: identify the inputs, the outputs, what each output element reads, and whether any thread needs information from another thread. It walks through several examples, including vector addition, scalar-vector multiplication, and ReLU, to show that map can involve one input array, multiple arrays, or a shared scalar, as long as each thread’s work remains independent. The CUDA implementation follows the same basic structure each time: compute a global thread index, check bounds, and apply one element-wise operation.

Finally, the chapter highlights two common mistakes that break the pattern or its safety: reading data that another thread is writing, which creates race conditions, and overwriting an input array when the original data is still needed. It emphasizes good habits such as using a separate output array, treating read-only values as normal kernel arguments, and remembering that threads do not run in a predictable order. By the end, the reader should be able to spot map problems quickly and write correct CUDA kernels for them with confidence.

The Map pattern: one thread, one output element, one independent transformation. Each thread reads from array A, applies function f, and writes the result to array B, with no thread depending on any other. The same pattern covers everything from simple vector addition to neural network activations like ReLU. One caution: ReLU’s max(0, A[i]) is an element-wise maximum (a map), whereas finding the single largest value across the whole array is a reduction (chapter 5), not a map.

The Parallel Checklist: four questions that confirm the Map pattern. Before writing a kernel, ask what your input and output are, what each output element needs, and what each thread needs from other threads. For the Map pattern, the answers are always the same: one or more input arrays in, one output array out, each output element reading only its same-index inputs, and no thread ever waiting on another.

Parallel Checklist for vector addition: all four boxes checked. Each thread reads A[i] and B[i], writes A[i]+B[i] to C[i], and needs nothing from any other thread: no dependency on any other output, no waiting. All four checklist answers confirm it: vector addition is a textbook Map pattern.

Parallel Checklist for scalar-vector multiplication: another clean Map. Each thread reads A[i] and the shared scalar k, writes k × A[i] to B[i], and needs nothing from any other thread: the scalar is read-only, so every thread can use it simultaneously without conflict.

Parallel Checklist for Vector ReLU. Each thread reads A[i], writes max(0, A[i]) to B[i], and needs nothing from any other thread: negative values get clamped to zero, while zero and positive values pass through unchanged, all in parallel.

• The map pattern applies the same operation to every element independently. No output depends on any other output: one thread per output element.
• The independence test: first, no output may depend on another output’s result (if one does, you need a different pattern); second, it’s a pure map only when every read lands at the thread’s own index i. A read at some other index is a gather, a close cousin that is still safe to run in parallel.
• Parallel Checklist: four questions (inputs, outputs, what each output reads, and what each thread needs from others) identify any map.
• Thread indexing: compute a unique global index with int idx = blockIdx.x * blockDim.x + threadIdx.x, then always check bounds with if (idx < n) to prevent out-of-bounds memory access.
• Scalars by value: pass simple values (like n or k) directly to kernels. Anything a kernel reads through a pointer must live in memory the device can reach, which is why arrays go through cudaMalloc and cudaMemcpy.
• Two mistakes to watch for: reading a neighbor while another thread overwrites it, and destroying an input by writing the result back into it.
• Map is everywhere, from scaling a vector by a constant to neural network activation functions like ReLU. Once you see the pattern, you’ll recognize it in NumPy, PyTorch, SQL, and beyond.

FAQ