GPU Challenges / CUDA

Image Blur Filter: CUDA Exercise

medium Runs on a real NVIDIA T4 GPU in your browser

Image blurring is a stencil computation: every output pixel depends on a small neighbourhood of input pixels. It is the pattern behind convolutions in image processing and CNNs, and it teaches 2D indexing and careful boundary handling.

The problem

Implement a simple 2D box blur filter for a 64x64 grayscale image. Each output pixel should be the average of itself and its 8 neighbors (3x3 kernel). Handle edge pixels by only averaging valid neighbors.

Implement a 3x3 box blur on a grayscale image while handling borders carefully so edge pixels only average valid neighbors.

Your goal

  • Translate 2D thread coordinates into image x and y positions.
  • Loop over the 3x3 neighborhood while ignoring out-of-bounds pixels.
  • Average the valid samples and write the blurred output pixel.

How to approach it

  1. Map each thread to one output pixel with a 2D grid.
  2. Loop over the 3x3 neighbourhood (offsets -1 to 1 in each direction).
  3. Skip neighbours outside the image, and count how many valid pixels you added.
  4. Write the sum divided by that count, so edge and corner pixels average only real neighbours.

Common mistakes

  • Dividing by 9 everywhere, which darkens the image border.
  • Reading out of bounds at the edges instead of skipping invalid neighbours.
  • Writing into the input buffer, so later threads read already-blurred values.

What success looks like

  • Center pixels should reflect a full 3x3 average.
  • Corner and edge pixels should not read outside the image.
  • Output values should remain in a valid unsigned char range.

Starter code

Start from this scaffold. Fill in the TODOs, then run it against the hidden tests.

image-blur.cu
1234567891011121314
#include <cuda_runtime.h>

// Kernel is launched as:  boxBlur<<<dim3(ceil(W/16),ceil(H/16)), dim3(16,16)>>>(input, output, width, height)
// gridDim  = (ceil(width/16), ceil(height/16))   blockDim = (16, 16)
// x (col) = blockIdx.x * blockDim.x + threadIdx.x
// y (row) = blockIdx.y * blockDim.y + threadIdx.y
__global__ void boxBlur(const unsigned char* input, unsigned char* output, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        // TODO: Sum all valid neighbors in 3x3 window (skip out-of-bounds)
        // TODO: Divide by neighbor count and write to output[y * width + x]
    }
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • 2D memory indexing
  • Boundary-safe stencil logic
  • Neighborhood accumulation and normalization

Learn them first:

More CUDA challenges