TensorTonic
Problems
Matrix Transpose as a CUDA Kernel
Linear AlgebraGPU Programming
Easy

Implement a CUDA kernel that computes the transpose of a 2D matrix. Given input matrix AAA of shape (M,N)(M, N)(M,N), write output BBB of shape (N,M)(N, M)(N,M) such that:

B[j,i]=A[i,j]B[j, i] = A[i, j]B[j,i]=A[i,j]

Both matrices are stored as row-major flat arrays on the device. AAA is indexed as A[i*N + j]\texttt{A[i*N + j]}A[i*N + j] and BBB is indexed as B[j*M + i]\texttt{B[j*M + i]}B[j*M + i], with 0≤i<M0 \le i < M0≤i<M and 0≤j<N0 \le j < N0≤j<N.

Loading visualization...

Examples

Input: M = 2, N = 3, A = [[1, 2, 3], [4, 5, 6]]

Output: B = [[1, 4], [2, 5], [3, 6]] (shape 3×23 \times 23×2)

Rows of AAA become columns of BBB.

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×33 \times 33×3)

A square matrix transposes in place of shape.

Hint 1

With a 2D launch, map threadIdx.y\texttt{threadIdx.y}threadIdx.y/blockIdx.y\texttt{blockIdx.y}blockIdx.y to the row iii of AAA and threadIdx.x\texttt{threadIdx.x}threadIdx.x/blockIdx.x\texttt{blockIdx.x}blockIdx.x to the column jjj of AAA, then bounds-check both before writing.

Hint 2

A naive transpose has coalesced reads on AAA but strided writes on BBB (or vice versa). That is expected for this version; a shared-memory tiled variant is not required here.

Requirements

  • Use one thread per output element with a 2D grid and 2D block layout
  • Treat A\texttt{A}A and B\texttt{B}B as row-major flat buffers using A[i*N + j]\texttt{A[i*N + j]}A[i*N + j] and B[j*M + i]\texttt{B[j*M + i]}B[j*M + i]
  • Bounds-check both axes so threads outside (M,N)(M, N)(M,N) skip their write
  • The solve\texttt{solve}solve signature must not change

Constraints

  • 1≤M,N≤40961 \le M, N \le 40961≤M,N≤4096
  • Input and output buffers are device pointers
  • Use a 16×1616 \times 1616×16 thread block
  • Values fit in 32-bit float
Try Similar Problems
Outer Product as a CUDA KernelEasyMatrix Addition as a CUDA KernelEasyHadamard Product as a CUDA KernelEasyRMS Normalization as a CUDA KernelMediumLayer Normalization as a CUDA KernelMediumMatrix Multiplication as a CUDA KernelMediumMatrix-Vector Multiplication as a CUDA KernelMediumVector Addition as a CUDA KernelEasyVector Subtraction as a CUDA KernelEasyL1 Normalization as a CUDA KernelMediumCosine Similarity as a CUDA KernelMediumCUDA Dot Product: Multi-Block ReductionMedium

Accepts: array

Accepts: number

Accepts: number

You must run your code first.

Kernel performance

Submit a correct solution to see your kernel performance.

Matrix Transpose as a CUDA Kernel

Matrix Transpose as a CUDA Kernel

Linear AlgebraGPU Programming
Easy

Implement a CUDA kernel that computes the transpose of a 2D matrix. Given input matrix AAA of shape (M,N)(M, N)(M,N), write output BBB of shape (N,M)(N, M)(N,M) such that:

B[j,i]=A[i,j]B[j, i] = A[i, j]B[j,i]=A[i,j]

Both matrices are stored as row-major flat arrays on the device. AAA is indexed as A[i*N + j]\texttt{A[i*N + j]}A[i*N + j] and BBB is indexed as B[j*M + i]\texttt{B[j*M + i]}B[j*M + i], with 0≤i<M0 \le i < M0≤i<M and 0≤j<N0 \le j < N0≤j<N.

Loading visualization...

Examples

Input: M = 2, N = 3, A = [[1, 2, 3], [4, 5, 6]]

Output: B = [[1, 4], [2, 5], [3, 6]] (shape 3×23 \times 23×2)

Rows of AAA become columns of BBB.

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×33 \times 33×3)

A square matrix transposes in place of shape.

Hint 1

With a 2D launch, map threadIdx.y\texttt{threadIdx.y}threadIdx.y/blockIdx.y\texttt{blockIdx.y}blockIdx.y to the row iii of AAA and threadIdx.x\texttt{threadIdx.x}threadIdx.x/blockIdx.x\texttt{blockIdx.x}blockIdx.x to the column jjj of AAA, then bounds-check both before writing.

Hint 2

A naive transpose has coalesced reads on AAA but strided writes on BBB (or vice versa). That is expected for this version; a shared-memory tiled variant is not required here.

Requirements

  • Use one thread per output element with a 2D grid and 2D block layout
  • Treat A\texttt{A}A and B\texttt{B}B as row-major flat buffers using A[i*N + j]\texttt{A[i*N + j]}A[i*N + j] and B[j*M + i]\texttt{B[j*M + i]}B[j*M + i]
  • Bounds-check both axes so threads outside (M,N)(M, N)(M,N) skip their write
  • The solve\texttt{solve}solve signature must not change

Constraints

  • 1≤M,N≤40961 \le M, N \le 40961≤M,N≤4096
  • Input and output buffers are device pointers
  • Use a 16×1616 \times 1616×16 thread block
  • Values fit in 32-bit float
Try Similar Problems
Outer Product as a CUDA KernelEasyMatrix Addition as a CUDA KernelEasyHadamard Product as a CUDA KernelEasyRMS Normalization as a CUDA KernelMediumLayer Normalization as a CUDA KernelMediumMatrix Multiplication as a CUDA KernelMediumMatrix-Vector Multiplication as a CUDA KernelMediumVector Addition as a CUDA KernelEasyVector Subtraction as a CUDA KernelEasyL1 Normalization as a CUDA KernelMediumCosine Similarity as a CUDA KernelMediumCUDA Dot Product: Multi-Block ReductionMedium

Accepts: array

Accepts: number

Accepts: number

You must run your code first.

Kernel performance

Submit a correct solution to see your kernel performance.

Problems
Loading...
1 / 1