CUDA Naive Matrix Transpose
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:
M = 64, N = 48
True
- The input values
M = 64andN = 48define the dimensions of the 2D matrixaas (M,N)=(64,48). - The
transpose_kernelfunction is launched with a 2D grid, where each thread reads one element ofaat position(i, j)and writes it to the swapped position(j, i)in the output matrixout. - The resulting
outmatrix has dimensions (N,M)=(48,64), which is the transpose of the original matrixa, i.e., out=aT. - The
runfunction compares the resultingoutmatrix with the reference transpose aT usingnp.allcloseand returnsTrueif 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
Background Knowledge
The problem involves using CUDA, a parallel computing platform and programming model developed by NVIDIA, to transpose a 2D matrix. In the context of GPU computing, a kernel is a small program that runs on the GPU, executing a specific task. The transpose operation is a fundamental linear algebra operation that swaps the rows and columns of a matrix. In this case, we're working with a 2D grid, where each thread reads one element of the input matrix a at position (i, j) and writes it to the corresponding position (j, i) in the output matrix out.
To understand the problem, it's essential to grasp the concept of memory coalescing, which refers to the ability of the GPU to access memory locations in a contiguous block, reducing memory access latency. In the context of matrix transposition, uncoalesced writes occur when threads write to non-contiguous memory locations, leading to reduced performance. The problem statement mentions that this "naive" transpose has uncoalesced writes, which will be addressed in a later problem using shared memory.
The CUDA memory model consists of several memory spaces, including global memory, shared memory, and register memory. Global memory is the largest memory space, but it's also the slowest. Shared memory, on the other hand, is a small, on-chip memory space that's shared among threads in a block, providing faster access times. Understanding the CUDA memory model and how to optimize memory access patterns is crucial for achieving high performance in GPU computing.
Continue the full explanation
You're reading the free preview. Unlock the complete walkthrough, the code editor, test runner and reference solution with Premium.
Editor locked
The code editor is locked for Pro problems. It is only available for free problems. Please upgrade to gain access to the code editor for all problems.