CUDA Shared-Memory Tiled Transpose

hardPROShared Memory & TilingCUDA Basics with Numba

Problem Statement

Transpose a matrix using a shared-memory tile so both the reads and the writes to global memory are coalesced: out = a.T. Assume M and N are multiples of the tile size TPB.

Background

The naive transpose has uncoalesced writes. Fix it by loading a TPB x TPB tile into shared memory with coalesced reads, cuda.syncthreads(), then writing the tile out to the transposed block position — reading tile[ty, tx] so the global write is again coalesced.

Your Task

Implement transpose_kernel and run(M=64, N=32) comparing to 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 = 32

Output

True

Reasoning

The input values M = 64 and N = 32 are used to allocate a matrix a of size $M \times N$ .
The transpose_kernel function is launched, which loads a $TPB \times TPB$ tile into shared memory with coalesced reads from the input matrix a.
The kernel then transposes the tile in shared memory and writes it out to the corresponding position in the output matrix out with coalesced writes, effectively computing out = a.T.
The resulting matrix out is compared to the reference solution a.T using np.allclose, and since the transpose operation is correctly implemented, the comparison returns True.

Constraints

Load tile[tx, ty] = a[x, y] (coalesced read), then cuda.syncthreads()
Write to the swapped block: out[blockIdx.yTPB + tx, blockIdx.xTPB + ty] = tile[ty, tx]
M and N are multiples of TPB

Learn this in

Moving Data: GPU Memory with NumbaCUDA with Numba Parallel Patterns: Reductions, Atomics & TilingCUDA with Numba ML InfrastructureSystems & Networking 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

Concept primerPREVIEW

Background Knowledge

The problem involves using CUDA (Compute Unified Device Architecture) to transpose a matrix. CUDA is a parallel computing platform and programming model developed by NVIDIA. It allows developers to harness the power of the GPU (Graphics Processing Unit) to perform general-purpose computing tasks. In the context of this problem, we're dealing with shared memory, which is a small amount of memory that is shared among threads in a block. Shared memory is much faster than global memory, but it's also limited in size.

To understand the problem, it's essential to know about coalesced memory access. In CUDA, coalesced memory access occurs when threads in a warp (a group of 32 threads) access memory locations that are contiguous and aligned. This type of access is much faster than uncoalesced access, where threads access non-contiguous or misaligned memory locations. The problem statement mentions that the naive transpose has uncoalesced writes, which means that the threads are writing to non-contiguous memory locations. To fix this, we need to use a shared-memory tile to load a block of data into shared memory with coalesced reads, and then write it out to the transposed block position with coalesced writes.

The concept of tiling is also crucial in this problem. Tiling involves dividing a large matrix into smaller blocks, called tiles, and processing each tile independently. This approach allows us to reduce the amount of memory accesses and improve the overall performance of the algorithm. In this case, we're using a TPB x TPB tile, where TPB is the number of threads per block. By loading a tile into shared memory and then writing it out to the transposed block position, we can ensure that both the reads and writes to global memory are coalesced.

PRO

Continue the full explanation

You're reading the free preview. Unlock the complete walkthrough, the code editor, test runner and reference solution with Premium.

PRO

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.

solution.py

Python 3.13

GPU · T4

Test results