CUDA Naive Matrix Transpose

mediumPROMemory & Data MovementCUDA Basics with Numba

Problem Statement

Transpose a 2D matrix with a 2D grid: out[j, i] = a[i, j], where a is (M, N) and out is (N, M).

Background

Each thread reads one element of a at (i, j) and writes it to the swapped position (j, i) in out. This "naive" transpose has uncoalesced writes — a later problem fixes that with shared memory — but it's the clearest place to start.

Your Task

Implement transpose_kernel and run(M=64, N=48) returning whether the result equals a.T.

How it is tested

Your solution must define a top-level function run(...) that allocates the inputs, copies them to the GPU, launches your @cuda.jit kernel, and returns a Python bool from np.allclose(gpu_result, reference). The grader prints run(...); the expected output is True.

Example

Input

M = 64, N = 48

Output

True

Reasoning

The input values M = 64 and N = 48 define the dimensions of the 2D matrix a as $(M, N) = (64, 48)$ .
The transpose_kernel function is launched with a 2D grid, where each thread reads one element of a at position (i, j) and writes it to the swapped position (j, i) in the output matrix out.
The resulting out matrix has dimensions $(N, M) = (48, 64)$ , which is the transpose of the original matrix a, i.e., $out = a^T$ .
The run function compares the resulting out matrix with the reference transpose $a^T$ using np.allclose and returns True if they are equal, which is the case for the given sample input.

Constraints

i, j = cuda.grid(2) index into a (shape M x N)
out has shape N x M; write out[j, i] = a[i, j]
Guard with if i < M and j < N

Learn this in

Moving Data: GPU Memory with NumbaCUDA with Numba GPU Architecture & the Memory HierarchyGPU Basics with Triton

More from CUDA Basics with Numba

CUDA Dot Producthard CUDA Atomic Max Reductionmedium CUDA Two-Stage Reduction with Atomicshard CUDA Atomic Histogrammedium CUDA Atomic Sum Reductionmedium