CUDA Shared-Memory Tiled Transpose
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:
M = 64, N = 32
True
- The input values
M = 64andN = 32are used to allocate a matrixaof size M×N. - The
transpose_kernelfunction is launched, which loads a TPB×TPB tile into shared memory with coalesced reads from the input matrixa. - The kernel then transposes the tile in shared memory and writes it out to the corresponding position in the output matrix
outwith coalesced writes, effectively computingout = a.T. - The resulting matrix
outis compared to the reference solutiona.Tusingnp.allclose, and since the transpose operation is correctly implemented, the comparison returnsTrue.
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
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.
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.