Implement a CUDA kernel that performs 2D max pooling on a single-channel input I of shape (H,W) with pool window (kH,kW) and stride (sH,sW). There is no padding.
For each output position (i,j):
output[i,j]=a∈[0,kH), b∈[0,kW)maxinput[i⋅sH+a, j⋅sW+b]The output has shape ((H−kH)/sH+1,(W−kW)/sW+1). Both 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 = 4, W = 4, kH = 2, kW = 2, sH = 2, sW = 2, input = [[1, 2, 3, 4], [5, 6, 7, 8], [9, 10, 11, 12], [13, 14, 15, 16]]
Output: [[6.0, 8.0], [14.0, 16.0]]
A standard non-overlapping 2×2 pool keeps the max of each disjoint window.
Input: H = 3, W = 3, kH = 2, kW = 2, sH = 1, sW = 1, input = [[1, 2, 3], [4, 5, 6], [7, 8, 9]]
Output: [[5.0, 6.0], [8.0, 9.0]]
With stride 1 the windows overlap, producing a 2×2 output of local maxima.
Each output cell owns an independent window of input reads. Initialise the reduction with -FLT_MAX (from <float.h>) so the first comparison always succeeds, then update with fmaxf.
Stride (sH,sW) controls how much the windows overlap or skip. The base index of the window owned by output (i,j) is (i*sH, j*sW); iterate a∈[0,kH), b∈[0,kW) over it.
Map blockIdx.y, threadIdx.y to the output row and blockIdx.x, threadIdx.x to the output column, and remember the output buffer has stride outW = (W - kW) / sW + 1.
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 2D max pooling on a single-channel input I of shape (H,W) with pool window (kH,kW) and stride (sH,sW). There is no padding.
For each output position (i,j):
output[i,j]=a∈[0,kH), b∈[0,kW)maxinput[i⋅sH+a, j⋅sW+b]The output has shape ((H−kH)/sH+1,(W−kW)/sW+1). Both 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 = 4, W = 4, kH = 2, kW = 2, sH = 2, sW = 2, input = [[1, 2, 3, 4], [5, 6, 7, 8], [9, 10, 11, 12], [13, 14, 15, 16]]
Output: [[6.0, 8.0], [14.0, 16.0]]
A standard non-overlapping 2×2 pool keeps the max of each disjoint window.
Input: H = 3, W = 3, kH = 2, kW = 2, sH = 1, sW = 1, input = [[1, 2, 3], [4, 5, 6], [7, 8, 9]]
Output: [[5.0, 6.0], [8.0, 9.0]]
With stride 1 the windows overlap, producing a 2×2 output of local maxima.
Each output cell owns an independent window of input reads. Initialise the reduction with -FLT_MAX (from <float.h>) so the first comparison always succeeds, then update with fmaxf.
Stride (sH,sW) controls how much the windows overlap or skip. The base index of the window owned by output (i,j) is (i*sH, j*sW); iterate a∈[0,kH), b∈[0,kW) over it.
Map blockIdx.y, threadIdx.y to the output row and blockIdx.x, threadIdx.x to the output column, and remember the output buffer has stride outW = (W - kW) / sW + 1.
Accepts: array
Accepts: number
Accepts: number
Accepts: number
Accepts: number
Submit a correct solution to see your kernel performance.