Gaussian Blur

原始题目：LeetGPU - Gaussian Blur

题目描述

编写一个 GPU 程序，对二维图像应用高斯模糊滤波器。给定一个以浮点数组表示的输入图像和一个高斯核，程序应计算图像与核的卷积。所有输入和输出均以行优先顺序存储。

高斯模糊通过对每个像素与其邻域的加权平均进行卷积来实现，权重由高斯核确定。对于输出像素位置 $(i, j)$：

$$output[i, j] = \sum_{m=-k_h/2}^{k_h/2}\sum_{n=-k_w/2}^{k_w/2} input[i+m, j+n] \times kernel[m+k_h/2, n+k_w/2]$$

其中 $k_h$ 和 $k_w$ 分别为核的高度和宽度。

实现要求

• 不允许使用外部库。
• solve 函数签名必须保持不变。
• 最终结果必须存储在 output 数组中。
• 边界条件使用零填充（将图像边界外的值视为零）。

示例

示例 1

Input:  image (5×5): 行优先 [[1,2,3,4,5],[6,7,8,9,10],[11,12,13,14,15],[16,17,18,19,20],[21,22,23,24,25]]
        kernel (3×3): [[0.0625, 0.125, 0.0625],[0.125, 0.25, 0.125],[0.0625, 0.125, 0.0625]]
Output: output (5×5): 第一行 [1.6875, 2.75, 3.5, 4.25, 3.5625], ...

示例 2

Input:  image (3×3): [[10,20,30],[40,50,60],[70,80,90]]
        kernel (3×3): [[0.1,0.1,0.1],[0.1,0.2,0.1],[0.1,0.1,0.1]]
Output: output (3×3): [[13,23,19],[31,50,39],[31,47,37]]

约束条件

• $1 \le input\_rows, input\_cols \le 4{,}096$。
• $3 \le kernel\_rows, kernel\_cols \le 21$，均为奇数。
• 所有核值非负且和为 $1.0$（已归一化）。
• 性能测试在 $input = 512 \times 512$，$kernel = 7 \times 7$ 的规模下进行。

解题思路

高斯模糊本质上是带零填充的二维卷积。与标准卷积的主要区别在于需要处理边界——越界的像素视为零。一种简单策略是让每个输出像素对应的线程检查输入坐标是否越界，越界则跳过。更高效的做法是预加载共享内存分块时包含 halo 区域，在边界处用零填充 halo。

代码实现

CUDA

#include <cuda_runtime.h>
__global__ void blur_kernel(const float* input, float* output, const float* kernel,
    int H, int W, int KH, int KW) {
    int col = blockIdx.x*blockDim.x+threadIdx.x, row = blockIdx.y*blockDim.y+threadIdx.y;
    if(row<H && col<W) {
        float sum=0.0f; int kh2=KH/2, kw2=KW/2;
        for(int ky=-kh2; ky<=kh2; ky++)
            for(int kx=-kw2; kx<=kw2; kx++) {
                int y=row+ky, x=col+kx;
                float v=(y>=0&&y<H&&x>=0&&x<W)?input[y*W+x]:0.0f;
                sum+=v*kernel[(ky+kh2)*KW+(kx+kw2)];
            }
        output[row*W+col]=sum;
    }
}
extern "C" void solve(const float* input, float* output, const float* kernel,
    int H, int W, int KH, int KW) {
    dim3 t(16,16), b((W+15)/16, (H+15)/16);
    blur_kernel<<<b,t>>>(input, output, kernel, H, W, KH, KW);
    cudaDeviceSynchronize();
}

Triton

import triton, triton.language as tl
@triton.jit
def blur_kernel(input_ptr, output_ptr, kernel_ptr, H, W, KH, KW, BLOCK: tl.constexpr):
    col = tl.program_id(0)*BLOCK + tl.arange(0, BLOCK)
    row = tl.program_id(1)*BLOCK + tl.arange(0, BLOCK)
    mask = (row[:,None]<H)&(col[None,:]<W)
    acc = tl.zeros((BLOCK,BLOCK), tl.float32)
    kh2,kw2=KH//2,KW//2
    for ky in range(-kh2, kh2+1):
        for kx in range(-kw2, kw2+1):
            y=row[:,None]+ky; x=col[None,:]+kx
            vm=(y>=0)&(y<H)&(x>=0)&(x<W)
            acc += tl.load(input_ptr+y*W+x, mask=vm, other=0.0) * tl.load(kernel_ptr+(ky+kh2)*KW+(kx+kw2))
    tl.store(output_ptr+row[:,None]*W+col[None,:], acc, mask=mask)