Solution: 1D Max Pooling

Problem: 1D Max Pooling

Submitted on: 7/12/2026, 11:38:03 PM

• Runtime: 442.05 μs
• Performance: 36.18 GFLOPS
• Tests: 6/6

#include <cuda_runtime.h>

constexpr size_t NUM_THREADS_PER_BLOCK = 256;

#define MAX(A, B) \
    (A > B ? A : B);

int ceil_div(int n1, int n2) {
    return ((n1 + n2 - 1) / n2);
}

struct Parameters {
    int kernel_size;
    int stride;
    int padding;
    int dilation;
    size_t H_in;
    size_t H_out;

    Parameters(int kernel_size, int stride, int padding, int dilation, size_t H_in) : 
        kernel_size(kernel_size), stride(stride), padding(padding), dilation(dilation), H_in(H_in)
    {
        H_out = ((H_in + 2 * padding - dilation * (kernel_size-1) - 1) / stride) + 1;

    }
};

__global__ void maxPoolingKernel(const float* input, float* output, Parameters* param) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < param->H_out) {
        size_t start_idx = param->stride * i - param->padding;
        float toBeMax = 0;
        if (start_idx < param->H_in) {
            toBeMax = input[start_idx];
        }
        for (int m = 0; m < param->kernel_size; ++m) {
            size_t idx = start_idx + param->dilation * m;
            if (idx < param->H_in) {
                toBeMax = MAX(toBeMax, input[idx]);
            }
        }
        output[i] = toBeMax;
    }
}

// Note: input, output are device pointers
extern "C" void solution(const float* input, int kernel_size, int stride, int padding, int dilation, float* output, size_t H) {
    Parameters param(kernel_size, stride, padding, dilation, H);

    Parameters* device_param;
    cudaMalloc(&device_param, sizeof(Parameters));
    cudaMemcpy(device_param, &param, sizeof(Parameters), cudaMemcpyHostToDevice);

    dim3 blockDim(NUM_THREADS_PER_BLOCK);
    dim3 gridDim(ceil_div(param.H_out, NUM_THREADS_PER_BLOCK));
    maxPoolingKernel<<<gridDim, blockDim>>>(input, output, device_param);
}