Implement a CUDA kernel that computes the transpose of a 2D matrix. Given input matrix A of shape (M,N), write output B of shape (N,M) such that:
B[j,i]=A[i,j]Both matrices are stored as row-major flat arrays on the device. A is indexed as A[i*N + j] and B is indexed as B[j*M + i], with 0≤i<M and 0≤j<N.
Input: M = 2, N = 3, A = [[1, 2, 3], [4, 5, 6]]
Output: B = [[1, 4], [2, 5], [3, 6]] (shape 3×2)
Rows of A become columns of B.
Input: M = 3, N = 3, A = [[1, 0, 2], [0, 3, 0], [4, 0, 5]]
Output: B = [[1, 0, 4], [0, 3, 0], [2, 0, 5]] (shape 3×3)
A square matrix transposes in place of shape.
With a 2D launch, map threadIdx.y/blockIdx.y to the row i of A and threadIdx.x/blockIdx.x to the column j of A, then bounds-check both before writing.
A naive transpose has coalesced reads on A but strided writes on B (or vice versa). That is expected for this version; a shared-memory tiled variant is not required here.
Accepts: array
Accepts: number
Accepts: number
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that computes the transpose of a 2D matrix. Given input matrix A of shape (M,N), write output B of shape (N,M) such that:
B[j,i]=A[i,j]Both matrices are stored as row-major flat arrays on the device. A is indexed as A[i*N + j] and B is indexed as B[j*M + i], with 0≤i<M and 0≤j<N.
Input: M = 2, N = 3, A = [[1, 2, 3], [4, 5, 6]]
Output: B = [[1, 4], [2, 5], [3, 6]] (shape 3×2)
Rows of A become columns of B.
Input: M = 3, N = 3, A = [[1, 0, 2], [0, 3, 0], [4, 0, 5]]
Output: B = [[1, 0, 4], [0, 3, 0], [2, 0, 5]] (shape 3×3)
A square matrix transposes in place of shape.
With a 2D launch, map threadIdx.y/blockIdx.y to the row i of A and threadIdx.x/blockIdx.x to the column j of A, then bounds-check both before writing.
A naive transpose has coalesced reads on A but strided writes on B (or vice versa). That is expected for this version; a shared-memory tiled variant is not required here.
Accepts: array
Accepts: number
Accepts: number
Submit a correct solution to see your kernel performance.