CUDA二维卷积算子从零完整实现详细教程与代码

本教程将指导您从零开始，使用CUDA构建一个高性能的2D卷积算子。即便您仅具备基础的CUDA编程知识，也能跟随本文的步骤，逐步实现一个在特定小尺寸下性能超越cudnn的卷积模块。内容涵盖卷积的基本概念，并深入探讨并行划分、数据迁移、预取等关键优化技术，最后提供完整代码和性能评测。

前言

当前深度学习飞速发展，CUDA作为底层加速的核心基础设施，掌握其编程技巧具有重要价值。本文将详细介绍如何利用CUDA实现一个2D卷积算子，实现流程相对简洁，最终版本在较小尺寸条件下能够获得比cudnn更优的性能。实际测试显示，在以下参数配置下，平均性能可达cudnn的1.2倍（该结果仅供参考，实际表现与cudnn配置相关，尺寸越小效果越明显）。

小提示： 测试时已排除cudnn初始化耗时，采用99轮平均时间。

const int inC = 6;
const int inH = 768;
const int inW = 512;
const int kernelH = 6;
const int kernelW = 6;
const int outC = 6;
const int outH = inH - kernelH + 1;
const int outW = inW - kernelW + 1

1. 卷积操作通俗介绍

1.1 数据布局 (data layout)

卷积计算主要针对图像进行，常见的RGB图像即包含三个通道的二维图像。数据通常以一维数组存储，并通过不同的布局方式索引。目前主流有NCHW和NHWC两种布局：

• n - batch size，即“图像”的个数
• c - channel num，通道数量
• h - height，每个通道的高度
• w - width，每个通道的宽度

NCHW布局按通道逐一读取整张图像，而NHWC则先读取所有通道同一位置的像素，再切换到下一位置。前者优先通道维度，后者优先空间位置。此外还有CHWN布局，但本文暂不涉及。

更多细节可参考英伟达官方文档：Developer Guide : NVIDIA Deep Learning cuDNN Documentation

NCHW layout

NHWC layout

本文按照NCHW数据格式实现算子。

1.2 直接卷积

最基本的直接卷积操作非常直观：想象一个滑动窗口（卷积核）在原矩阵上移动，每次与对应位置的元素进行点积并求和，结果填入输出矩阵的相应位置。下图能清晰展示这一过程（常被称为“对对碰”）。

图源：国科大模式识别课程

对于多输入通道的情况，分别对每个通道相同位置与卷积核进行“对对碰”，并将各通道结果累加作为输出值；当存在多输入多输出通道时，只需对每个输出通道重复上述操作。

有关通道的深入理解，推荐参考文章：CNN卷积核与通道讲解

注意：本文不涉及padding和stride的实现。

padding：对图像边缘进行填充，避免边缘特征被忽略，同时保持特征图尺寸不缩小。

stride：窗口滑动的步长，默认为1。例如stride=2时，下一次卷积的中心会跳过1个元素。

1.3 其他卷积计算方法

除了直接卷积，还有以下几种常见方法，读者可自行了解：

• Img2col：将图像展开为行向量，卷积核展开为列向量，转化为矩阵乘法。
• FFT method：利用傅里叶变换在频域进行卷积，计算量更小。
• Winograd Algorithm：将图像变换到另一个空间，减少乘法运算次数。

2. 整体实现思路

2.1 block与thread划分

我们需要考虑如何对多通道图像矩阵进行block和thread的划分。不同的切分方式会影响block在SM上的流转效率。本文采用 8×4的block，共32个线程（一个warp的大小），这样每个block上到SM时无需与其他block拼接线程，避免额外开销。

每个线程负责一个“位置”的计算，而不是单个元素。以16×20的矩阵为例，划分结果如下图（使用笛卡尔坐标系，与行主序不同）。

2.2 数据转移

关于位置和规模(size)

为什么说一个block有32个“位置”而不是32个元素？以3×3卷积核为例，卷积操作遍历原矩阵所有元素，但按中心点位置计数，结果如下：

注意：这里仅示意卷积中心点的范围。若按左上角元素来看，范围是：

参数解释：

• row_num：原矩阵一行元素数目
• inH, inW：原矩阵的高度和宽度
• kernelH, kernelW：卷积核的高度和宽度
• outH, outW：输出矩阵的高度和宽度

输出矩阵的规模为：

请注意：大小和位置下标有区别，一个从1开始数，一个从0开始数。

一个block的数据转移

对输出矩阵进行block划分更合理，因为32个位置恰好对应输出矩阵的32个位置。一个block需要读取的原矩阵范围是：

22×18的输入产生20×16的输出。以(0,0)和(4,1)的block为例，涉及原矩阵范围如下图：

蓝色区域为一个block需要用到的原矩阵元素。8×4的block需要读取 10×6 的元素（即 +kernelH-1）。输出矩阵则“一个萝卜一个坑”，无需额外考虑。

2.3 计算逻辑

不考虑channel（单输入单输出单卷积核）：只做三件事：

1. 将block对应的数据转移到SMEM中
2. 利用线程tid计算对应输出位置的结果
3. 将结果写回输出矩阵

只考虑inC：加一层循环，每个线程计算多个in channel的数据并累加，需要用寄存器存储中间结果。

考虑inC与outC：开大空间，让线程存储多个outC的中间结果，分别累加，最后分别写回。

3. 详细实现过程

3.1 整体实现思路

先从最简单的CPU串行版本开始，保证正确性，再逐步进行并行化改造。直接卷积运算过程相对简单，无需参考太多资料。

3.1.1 CPU串行版本的卷积算子

#define element_type float
#define OFFSET(row, col, ld) ((row) * (ld) + (col))
/*
    @brief: 串行卷积实现 CPU代码 NCHW
    @param in inC inH inW: 输入矩阵(数组) channel height width
    @param out outC outH outW: 输出矩阵 channel height width
    @param kernel kernelH kernelW: 卷积核 height width
*/
void serial_convolution(element_type *in, element_type *out, element_type *kernel, int batch_size,
                        int inC, int inH, int inW,
                        int outC, int outH, int outW,
                        int kernelH, int kernelW)
{
    float val;
    int out_pos, in_pos, kernel_pos;
    for (int oc = 0; oc < outC; oc++) // 每个输出通道
    {
        // 对一个位置的操作 用当前输入channel卷积去对相应的输出channel
        // 保证每个outChannel都是多inChannel累积的结果
        for (int i = 0; i < outH; i++)
        {
            for (int j = 0; j < outW; j++)
            {
                val = 0; // 避免累积和需要多次读取写入
                out_pos = oc * outH * outW + OFFSET(i, j, outW);
                for (int ic = 0; ic < inC; ic++) // 对每个输入通道
                {
                    for (int ii = 0; ii < kernelH; ii++)
                    {
                        for (int jj = 0; jj < kernelW; jj++)
                        {
                            in_pos = ic * inH * inW + OFFSET(i + ii, j + jj, inW);
                            kernel_pos = oc * kernelH * kernelW + OFFSET(ii, jj, kernelW);
                            val += in[in_pos] * kernel[kernel_pos];
                        }
                    }
                }
                out[out_pos] = val; // 与cudnn计算结果为相反数
            }
        }
    }
}

这是最朴素的O(n⁵)算法，但仅关注正确性，CUDA上不会运行这一版本。重点在于尽量减少访存，多用本地存储（寄存器、SMEM）代替GMEM访问。

3.1.2 循环顺序设计

循环嵌套顺序：outC → H → W → inC → kernelH → kernelW。设计理由：

• outC在最外层：一次循环内完成一个输出通道的所有计算，利用NCHW布局的局部性。
• H先于W：行主序存储，先H更高效。
• inC放在H/W之后：避免为每个输出位置维护多个中间变量，只需一个float即可。

小提示： 将inC放在H/W之后，可显著减少寄存器使用量。

3.2 线程任务均分

block大小为8×4（32线程），但in矩阵数据远多于32个元素，需要公平分配任务。本文采用“均分”思想，让每个线程处理差不多数量的元素。

方案①：边缘线程负责白色区域（效率低，不采用）

方案②：平均划分（推荐）

定义变量：

// 分块边界 boundary是限制正常范围 edge是需要补的范围
int row_boundary = outH / BLOCK_HEIGHT - 1,
    col_boundary = outW / BLOCK_WIDTH - 1;
int row_edge = outH % BLOCK_HEIGHT, col_edge = outW % BLOCK_WIDTH;
...
int single_trans_ele_num = 4;                               // 线程一次转移的数据数
int cur_in_block_height = BLOCK_HEIGHT + KERNEL_HEIGHT - 1, // 读入in的block height
    cur_in_block_width = BLOCK_WIDTH + KERNEL_WIDTH - 1,    // 读入in的block width
    in_tile_thread_per_row,                                 // 以tile为单位转移数据，一行需要的thread数
    in_tile_row_start,                                      // tile的行起始位置
    in_tile_col,                                            // tile的列
    in_tile_row_stride;                                     // tile行跨度
// 修正边缘block尺寸
if (block_row == row_boundary){
    cur_in_block_height = BLOCK_HEIGHT + row_edge + kernelH - 1;
}
if (block_col == col_boundary){
    cur_in_block_width = BLOCK_WIDTH + col_edge + kernelW - 1;
}
in_tile_thread_per_row = cur_in_block_width / single_trans_ele_num;
in_tile_row_start = tid / in_tile_thread_per_row;
in_tile_col = tid % in_tile_thread_per_row * single_trans_ele_num;
in_tile_row_stride = thread_num_per_block / in_tile_thread_per_row;

3.2.1 “block”设计与修正

由于out矩阵的尺寸可能不是block的整数倍，需要修正边缘block的大小。例如16×20的out矩阵，block网格只有16×20，但实际out有18×22，需要扩大边缘block覆盖全部范围。

3.2.2 线程行为指定

通过tile方式分配任务，每个线程负责连续的4个元素（利用float4向量指令，single_trans_ele_num=4）。图示：

循环结构：

for (int i = 0; i < cur_in_block_height && in_tile_row_start < cur_in_block_height;
        i += in_tile_row_stride){
    /*do something*/
}

注意： 由于线程数32可能不是每行tile数的整数倍，会出现某些线程处理多个行的情况，但行为一致，不影响结果。

3.3 预取机制

计算效率远高于访存，因此需要预取数据。整体数据传输逻辑：GMEM → SMEM → register → GMEM。伪代码框架：

初始化所有变量并修正block规模;
分配shared memory;
// 预读取第一个channel的数据
for (int i = 0; i < cur_in_block_height && in_tile_row_start < cur_in_block_height;
        i += in_tile_row_stride){
    把in中的数据从GMEM转到SMEM;
}
// 预读取第一个kernel的数据
if (thread_row >= 0 && thread_row < KERNEL_HEIGHT && thread_col == 0){
    把kernel的数据从GMEM转到SMEM;
}
__syncthreads();
// 先oc再ic的循环
for (int oc = 0; oc < outC; oc++){
    for (int ic = 0; ic < inC; ic++){
        // 计算当前ic与oc的结果，存到register
        ...
        // 读取下一个in channel数据
        if (ic + 1 < inC){
            for (...){
                读取下一个channel的数据;
            }
        }
        __syncthreads();
    }
    // 写回当前outC的数据
    for (...){
        写回;
    }
    // 预读取下一个kernel数据
    if (oc + 1 < outC){
        读取下一个kernel;
    }
    // 预读取第一个in channel数据（从头开始）
    for (...){
        读取第一个channel的数据;
    }
}

3.4 一些杂项细节

3.4.1 中间结果存储设计

先oc后ic的循环顺序，寄存器使用量约为oc×2（每个线程需维护多个输出位置）。为避免寄存器溢出到local memory（位于GMEM），尽量控制寄存器数量。本文直接开4倍空间保证安全。

3.4.2 下标计算

利用向量指令读取4个数据：

FETCH_FLOAT4(load_reg[0]) =
    FETCH_FLOAT4(in[begin_pos + OFFSET(in_tile_row_start + i, in_tile_col, inW)]);
s_in[in_tile_row_start + i][in_tile_col] = load_reg[0];
s_in[in_tile_row_start + i][in_tile_col + 1] = load_reg[1];
s_in[in_tile_row_start + i][in_tile_col + 2] = load_reg[2];
s_in[in_tile_row_start + i][in_tile_col + 3] = load_reg[3];

计算中间结果位置：

temp_pos = i / out_tile_row_stride + j +
           oc * (cur_out_block_height / out_tile_row_stride + 1);

写出结果：

out_pos = oc * outH * outW +
          block_row * BLOCK_HEIGHT * outW + block_col * BLOCK_WIDTH +
          OFFSET(out_tile_row_start + i, out_tile_col + j, outW);

4. 性能测试

虽然为娱乐测试，但结果仍具参考价值。代码性能受 channel 数目影响较大，不同数据规模下性能在 cudnn 的 1/10 到 10 倍之间波动。建议读者自行测试并进一步优化。

常见问题：

• Q：为什么我的结果与cudnn相反数？
  A：串行版本中代码注释已说明，这是因为计算顺序或符号问题，不影响性能对比，可自行调整符号。
• Q：边缘block修正后，线程任务分配是否仍然均匀？
  A：边缘block尺寸变化会导致部分线程承担更多任务，但整体性能影响较小，可进一步优化。
• Q：为什么必须使用float4向量指令？
  A：一次性读取4个float数据，充分利用显存带宽，提升数据转移效率。

结语

本教程从零开始，逐步构建了一个基于CUDA的2D卷积算子。通过理解数据布局、block划分、数据转移和预取机制，您可以在小尺寸参数下获得优于cudnn的性能。虽然代码仍有优化空间（如边缘处理、公共表达式提取），但它为深入学习CUDA算子优化提供了坚实基础。鼓励读者在此基础上继续探索更高效的卷积实现。

你是一名 AI 行业编辑，请围绕下面这条热点输出一份资讯解读：
热点：CUDA二维卷积算子从零完整实现详细教程与代码要求：
1. 先用一句话解释这条热点在讲什么
2. 再总结它为什么重要
3. 说明会影响哪些 AI 产品或内容方向
4. 最后给出 3 个适合资讯站使用的标题