One operation, many numbers

Idea 1 of 5

One calculation for five inputs

The PyTorch operation multiplies every input by the same amount. Follow the five numbers through the picture, then try a different multiplier.

y = x * 2

Browser simulation · this example has its own controls

Predict the outputs, then press Run.

Idea 2 of 5

Which number does each thread use?

The previous picture showed the same calculation applied to five inputs. A GPU can divide that work among many executions of a small program: each reads an input, multiplies it, and writes a result. These executions are called threads, and many can make progress in parallel.

In this picture, one thread handles one input. Its index tells it which input to read and where to write the result. This is a simplified mapping; PyTorch chooses the actual implementation.

y = x * 2

Browser simulation · this example has its own controls

Predict the outputs, then press Run.

Select a thread to follow its input and output.

Idea 3 of 5

Threads are organised into blocks

How do we ask the GPU to start those threads? We specify groups rather than starting each thread separately. The GPU can distribute these groups, and threads within a group can cooperate when an algorithm needs them to.

Here each thread multiplies its own input independently, so it needs no cooperation. Threads in a block do not necessarily all execute at once. Later, we will use cooperation to add many inputs into one result.

y = x * 2

For our example we choose four threads per block. Two blocks provide eight threads for five inputs. Each thread has a position within its block. Its index across the array is: block number × threads per block + position within the block.

Block 1, position 0 gives 1 × 4 + 0 = 4: the fifth input. The extra three threads still exist, but have no input. A check makes them skip the read, multiplication and write. Their cards are grey and say “No calculation”; select one to inspect its index. The next experiment lets you turn that check off.

Browser simulation · this example has its own controls

Predict the outputs, then press Run.

Select a thread to follow its input and output.

Idea 4 of 5

More threads than numbers

Keep the same five inputs and two blocks of four threads. The extra three threads must skip the arrays because their indices have no input or output element.

y = x * 2

Here, the bounds check asks whether a thread’s index is smaller than the number of inputs. Indices 0–4 pass; indices 5–7 skip the read, multiplication and write. Turn the check off below to see why it matters.

Browser simulation · this example has its own controls

Predict the outputs, then press Run.

Select a thread to follow its input and output.

Idea 5 of 5

From Python to a GPU kernel

We have followed the work in pictures. Now we can write the small program each thread follows and show the code that starts the groups.

Each thread executes the complete function below with its own index. Here is a CUDA kernel and the launch that starts it, putting together the index, calculation and bounds check you have just explored.

y = x * 2

PyTorch selects an optimised implementation and handles its launch; y = x * 2 does not choose blocks or threads. CUDA lets us write the GPU function and choose those settings ourselves. The compiler also optimises our code. Our runKernel is a separate teaching function; PyTorch does not call it when evaluating this expression. The example makes the work visible and is not intended to outperform PyTorch.

__global__ void runKernel(const float* input, float* output, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        output[i] = input[i] * 2.0f;
    }
}

runKernel<<<2, 4>>>(d_input, d_output, 5);

1. Find the input. blockIdx.x is the block number, blockDim.x is its size, and threadIdx.x is the position within it. Their calculation gives the integer index i. The .x selects the first dimension of our one-dimensional launch.
2. Check the index. n is the input count. if (i < n) lets only existing elements reach the calculation.
3. Calculate one result. The assignment reads input[i], multiplies it and writes output[i]. 2.0f is a floating-point number; ; ends the statement.

The launch’s <<<2, 4>>> starts two blocks of four threads. d_input and d_output refer to arrays already allocated in GPU memory; the final 5 supplies n. The download below includes that allocation and the copies.

Browser simulation · this example has its own controls

Predict the outputs, then press Run.

Select a thread to follow its input and output.

nvcc -std=c++17 hello.cu -o hello
./hello