一维网格:最基础的配置

说实话,一维网格是我用得最多的配置。为什么?因为大部分数据处理任务,本质上都是一维的。你想想看,一个数组有100万个元素,每个元素做一次计算,这不就是一维问题吗?

声明一维网格很简单:

dim3 grid(256);    // 等价于 dim3 grid(256, 1, 1)
dim3 block(256);   // 等价于 dim3 block(256, 1, 1)

kernel<<<grid, block>>>(d_data, N);

这里有个细节我要强调一下。dim3是个三维向量,但你只传一个参数时,后面两个维度默认是1。我个人习惯是,即使只用一维,也明确写成dim3 grid(256, 1, 1),这样代码读起来更清晰。

核心要点:一维网格适合处理线性数据,比如数组、向量、列表。每个线程处理一个或连续几个元素。

什么时候用一维网格?

我在项目中遇到过这样一个场景:需要对一个100万长度的浮点数组做逐元素归一化。这就是典型的一维问题。每个线程处理一个元素,网格大小就是ceil(N / blockSize)

常见的适用场景:

  • 向量加法、点乘、标量运算
  • 数组排序、归约、扫描
  • 图像的一维滤波(比如水平方向的高斯模糊)
  • 任何数据量很大但维度单一的任务

小技巧:计算网格大小时,记得处理边界。我一般这样写:

int gridSize = (N + blockSize - 1) / blockSize;
kernel<<<gridSize, blockSize>>>(d_data, N);

然后在核函数里加个if (idx < N)判断,防止越界。

二维网格:处理图像和矩阵

二维网格就更有意思了。我第一次用二维网格是在做图像卷积的时候。一张1920x1080的图片,每个像素需要访问周围3x3的邻域。如果用一维网格,你得手动计算行号和列号,代码写起来很别扭。

声明二维网格:

dim3 block(16, 16);      // 16x16 的线程块
dim3 grid(
    (width + block.x - 1) / block.x,
    (height + block.y - 1) / block.y
);

kernel<<<grid, block>>>(d_image, width, height);

你看,这样每个线程块天然对应图像的一个小区域。在核函数里,通过blockIdx.xblockIdx.y就能直接定位到像素坐标,多直观。

为什么用二维? 因为数据本身是二维的。用二维网格,你的代码逻辑和数据结构完全对应,可读性大大提升。而且,GPU的线程调度器对二维网格也有优化,访存模式更友好。

二维网格的典型场景

嗯,这里要注意,不是所有二维数据都适合用二维网格。我见过有人把一维数组硬套成二维网格,结果性能反而下降了。什么时候该用?

  • 图像处理:滤波、边缘检测、色彩空间转换
  • 矩阵运算:矩阵乘法、转置、LU分解
  • 物理模拟:二维网格上的有限差分、流体模拟
  • 机器学习:卷积层的前向传播、池化操作

避坑指南:我曾经在做一个矩阵转置时,用了32x32的线程块,结果共享内存爆了。后来才发现,每个线程块需要的共享内存是block.x * block.y * sizeof(float),32x32就是4KB,再算上其他变量,很容易超限。建议二维线程块大小控制在16x16以内,除非你很清楚硬件限制。

三维网格:立体数据的处理

三维网格用得相对少一些,但一旦需要,它就是不可替代的。我记得有一次做CT图像重建,数据是512x512x512的体素。用一维或二维网格处理,代码复杂度会爆炸。三维网格直接搞定。

声明三维网格:

dim3 block(8, 8, 8);     // 三维线程块
dim3 grid(
    (depth + block.z - 1) / block.z,
    (height + block.y - 1) / block.y,
    (width + block.x - 1) / block.x
);

kernel<<<grid, block>>>(d_volume, width, height, depth);

注意看,三维网格的声明方式和一维、二维完全一致,只是多了一个维度。在核函数里,你可以用blockIdx.xblockIdx.yblockIdx.z来定位体素坐标。

关键限制:三维网格的每个维度最大是65535(计算能力3.0以上是2^31-1)。但实际中,你很少会用到这么大的网格,因为线程块数量太多反而会降低占用率。

三维网格的使用场景

说白了,三维网格就是为三维数据量身定做的。我总结了几类典型应用:

  • 体数据处理:CT/MRI图像重建、地震数据体分析
  • 三维物理模拟:流体动力学、电磁场计算、分子动力学
  • 深度学习:3D卷积神经网络、点云处理
  • 科学计算:三维有限元分析、气候模型

个人经验:三维网格的线程块大小,我一般选4x4x4或8x8x8。为什么?因为每个线程块的总线程数是block.x * block.y * block.z,8x8x8=512,刚好是warp大小的整数倍。而且三维网格的访存模式比较复杂,线程块太大容易导致缓存抖动。

三种网格的对比与选择

你可能会问:到底该用哪种网格?我直接给你一个决策树:

数据类型 推荐网格 理由
一维数组(向量) 一维网格 代码简单,访存连续,性能最优
二维数组(图像、矩阵) 二维网格 坐标映射直观,共享内存利用方便
三维数组(体数据) 三维网格 天然匹配数据结构,代码可读性强
一维数据但需要二维索引 一维网格 手动计算索引,避免网格维度浪费

重要提醒:不要为了用三维而用三维。我见过有人把一维数组包装成三维网格,结果每个维度只有几个线程块,GPU的并行度完全没发挥出来。记住,网格的每个维度至少要有几十个线程块,才能充分利用SM的并行能力。

实际项目中的选择策略

最后,我分享一个实际项目中的决策方法。在做性能调优时,我会先问自己三个问题:

  1. 数据是几维的? 一维用一维网格,二维用二维网格,三维用三维网格。这是最自然的选择。
  2. 线程块大小是否合理? 每个线程块至少128个线程,最多1024个。我一般选256或512。
  3. 网格大小是否足够? 网格的每个维度至少几十个线程块,否则并行度不够。

如果这三个问题都答不上来,那就从一维网格开始试。一维网格的调优最简单,性能也最容易预测。等代码跑通了,再考虑是否要改成二维或三维。

总结一句话:网格的维度选择,本质上是在「代码可读性」和「性能」之间做权衡。对于大多数场景,选择与数据结构匹配的网格维度,就是最好的选择。

好了,这一章的内容就到这里。下一章我们会深入线程块的配置,包括线程块大小对性能的影响,以及如何选择最优的线程块尺寸。到时候我会分享一些具体的调优案例,都是我在项目中踩过的坑。