Implement a CUDA kernel that performs a 2D cross-correlation of a single-channel input I of shape (H,W) with a single-channel kernel K of shape (kH,kW). This is the operation deep learning frameworks call "convolution" (no kernel flip). Stride is 1 and there is no padding (valid).
For each output position (i,j):
output[i,j]=a=0∑kH−1b=0∑kW−1input[i+a,j+b]⋅kernel[a,b]The output has shape (H−kH+1,W−kW+1). All three buffers are stored as row-major flat arrays, so element (i,j) of an R×C matrix lives at linear index i⋅C+j.
Input: H = 3, W = 3, kH = 2, kW = 2, input = [[1, 2, 3], [4, 5, 6], [7, 8, 9]], kernel = [[1, 0], [0, -1]]
Output: [[-4.0, -4.0], [-4.0, -4.0]]
Each 2×2 window contributes top-left−bottom-right, which is −4 everywhere on this diagonal ramp.
Input: H = 2, W = 3, kH = 1, kW = 2, input = [[1, 2, 3], [4, 5, 6]], kernel = [[1, 1]]
Output: [[3.0, 5.0], [9.0, 11.0]]
A horizontal sum of every adjacent pair in each row.
Each output position owns an independent kH×kW window of input reads, so threads do not need to synchronize or share state.
Map blockIdx.y, threadIdx.y to the output row i and blockIdx.x, threadIdx.x to the output column j, and remember the output buffer has its own stride outW=W−kW+1.
For small kernel sizes the inner double loop is tight enough that the compiler unrolls it; accumulating into a local float avoids repeated global memory writes.
Accepts: array
Accepts: array
Accepts: number
Accepts: number
Accepts: number
Accepts: number
Submit a correct solution to see your kernel performance.
Implement a CUDA kernel that performs a 2D cross-correlation of a single-channel input I of shape (H,W) with a single-channel kernel K of shape (kH,kW). This is the operation deep learning frameworks call "convolution" (no kernel flip). Stride is 1 and there is no padding (valid).
For each output position (i,j):
output[i,j]=a=0∑kH−1b=0∑kW−1input[i+a,j+b]⋅kernel[a,b]The output has shape (H−kH+1,W−kW+1). All three buffers are stored as row-major flat arrays, so element (i,j) of an R×C matrix lives at linear index i⋅C+j.
Input: H = 3, W = 3, kH = 2, kW = 2, input = [[1, 2, 3], [4, 5, 6], [7, 8, 9]], kernel = [[1, 0], [0, -1]]
Output: [[-4.0, -4.0], [-4.0, -4.0]]
Each 2×2 window contributes top-left−bottom-right, which is −4 everywhere on this diagonal ramp.
Input: H = 2, W = 3, kH = 1, kW = 2, input = [[1, 2, 3], [4, 5, 6]], kernel = [[1, 1]]
Output: [[3.0, 5.0], [9.0, 11.0]]
A horizontal sum of every adjacent pair in each row.
Each output position owns an independent kH×kW window of input reads, so threads do not need to synchronize or share state.
Map blockIdx.y, threadIdx.y to the output row i and blockIdx.x, threadIdx.x to the output column j, and remember the output buffer has its own stride outW=W−kW+1.
For small kernel sizes the inner double loop is tight enough that the compiler unrolls it; accumulating into a local float avoids repeated global memory writes.
Accepts: array
Accepts: array
Accepts: number
Accepts: number
Accepts: number
Accepts: number
Submit a correct solution to see your kernel performance.