Matrix Multiplication 2: First CUDA Kernel

You have a real your selected GPU. Take the same N = 4096 matrix multiplication you built on one CPU core and launch it across thousands of CUDA threads. Each thread owns one output element of C. Write the kernel, run it on the GPU, and watch your first CUDA computation come to life.

Languages: CUDA C++, Python, Rust, HIP C++

Lesson

This is the jump from one CPU core to a real GPU. The mathematics has not changed: C[row, col] is still the dot product of a row of A and a column of B. What changed is the mapping. On the CPU, one loop eventually visits every output element. On the GPU, a two-dimensional grid gives each thread one row and one column. The runner moves A and B to your selected GPU, launches your kernel, moves C back, checks hidden output positions, and records both kernel-only and end-to-end time.

This stage assumes square matrices. A general implementation must also handle non-square dimensions and tile tails; do not infer that coverage from these tests.

What you will learn

Map a two-dimensional CUDA thread index to one output element of C.
Reuse the row-major indexing rule from the CPU baseline inside a CUDA kernel.
Distinguish the kernel's execution time from the full H2D, kernel, and D2H trip.

Starter code

    // One thread computes one output element C[row, col].
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            // TODO: multiply A[row, k] by B[k, col] and add it to sum.
        }
        // TODO: write sum to C[row, col].
    }