CUDA学习之路[4]——CUDA全局坐标计算

2706 字
14 分钟
CUDA学习之路[4]——CUDA全局坐标计算

目录#

引言#

在上一篇中,我们理清了 GPU 的硬件架构与 CUDA 的线程模型:Grid → Block → Thread

但很多同学一到写代码就卡住,最先卡住的就是算不对全局索引

int i = blockIdx.x * blockDim.x + threadIdx.x;
本篇学习目标
  • 彻底掌握一维、二维、三维全局索引的计算公式
  • 理解 gridDimblockDim 与启动配置之间的对应关系
  • 学会根据数据形状选择合适的网格/块维度

思考题:假如你有一个 1000×1000 的图像要处理,你是开一个 1000×1000 的二维线程块,还是开一个一维的 1,000,000 线程?为什么?答案将在文中揭晓。

核心概念速查#

CUDA 为每个线程提供了四个内置变量,用于定位自己在整个任务中的位置:

变量含义维度范围
threadIdx.x/y/z线程在块内的局部索引0 ~ blockDim-1
blockIdx.x/y/z线程块在网格内的索引0 ~ gridDim-1
blockDim.x/y/z每个块每维的线程数由启动参数 <<<..., threads>>> 决定
gridDim.x/y/z网格每维的块数由启动参数 <<<grid, ...>>> 决定

全局索引的本质跳过前面所有块的线程,再加上我在当前块内的偏移

全局坐标计算方式#

无论数据是几维,内存中都是一维线性排布。全局坐标计算的核心任务,就是将多维逻辑索引映射为一维物理地址

2. 一维坐标计算#

公式#

int global_id = blockIdx.x * blockDim.x + threadIdx.x;

图示#

一维全局索引计算方式
一维全局索引计算方式

  • blockIdx.x * blockDim.x:计算当前块之前一共有多少个线程(基地址)。
  • + threadIdx.x:加上当前线程在块内的偏移。

边界保护#

if (global_id < N) {
// 安全处理
}
为什么必须加这个判断?

:网格覆盖的线程总数往往是块大小的整数倍,可能超过数据总量 N。不加判断会越界访问。

2.2 二维坐标计算#

二维数据(如图像)需要两个全局索引:rowcol

公式#

// 列方向,当前块之前的数量 * 每块宽度 + 当前块内的线程偏移
int col = blockIdx.x * blockDim.x + threadIdx.x;
// 行方向,当前块之前的数量 * 每块高度 + 当前块内的线程偏移
int row = blockIdx.y * blockDim.y + threadIdx.y;
// 我们只需要利用行优先规则转换
if (row < height && col < width) {
int global_idx = row * width + col; // 转为一维线性地址
}

图示#

二维全局索引计算方式
二维全局索引计算方式

RGB图像转灰度图#

我们来看一下如何使用CUDA来将彩色图片转为灰度图。

安装opencv
Terminal window
sudo apt update
sudo apt install libopencv-dev

在一个Kernel中,每个线程负责处理图像中的一个像素点,我们需要利用上述公式计算出每个像素的行和列。

// 输入图像在显存中是按 uchar3 (RGB三通道) 紧密排列的
__global__ void rgbToGray(const uchar3* d_img, unsigned char* d_gray, 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) {
// 二维转一维索引
int idx = y * width + x;
// 读取原始 RGB 像素
uchar3 pixel = d_img[idx];
// 计算灰度值 (固定加权公式)
// Y = 0.299R + 0.587G + 0.114B
unsigned char gray = static_cast<unsigned char>(0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z);
// 写回显存
d_gray[idx] = gray;
}
}

在主机端,我们根据图像实际尺寸来配置Grid和Block,确保能够覆盖所有的像素。

int main() {
// ... 读取图像,获取 width 和 height ...
// 定义线程块大小:16x16 是常见且高效的选择
dim3 blockSize(16, 16); // dim3是CUDA中的数据结构
// 计算网格大小:用图像尺寸除以块大小并向上取整
// 公式:(N + M - 1) / M
dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
(height + blockSize.y - 1) / blockSize.y);
// 启动 Kernel
rgbToGray<<<gridSize, blockSize>>>(d_img, d_gray, width, height);
// ... 后续同步与结果保存 ...
}

完整的代码如下:

#include <iostream>
#include <opencv2/opencv.hpp>
#include <cuda_runtime.h>
#include <string>
__global__ void rgbToGray(const uchar3* d_img, unsigned char* d_gray, 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) {
int idx = y * width + x;
uchar3 pixel = d_img[idx];
// 使用加权平均法将RGB转换为灰度
d_gray[idx] = static_cast<unsigned char>(0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z);
}
}
using namespace std;
int main() {
string imagePath = "input.png";
cv::Mat img = cv::imread(imagePath);
// 检查图像是否加载成功
if (img.empty()) {
cerr << "无法加载图像: " << imagePath << endl;
return -1;
}
int width = img.cols;
int height = img.rows;
int channels = img.channels();
uchar* d_img; // GPU上的原始图像数据
unsigned char* d_gray; // GPU上的灰度图像
cout << "图像尺寸: " << width << "x" << height << ", 通道数: " << channels << endl;
// 定义CUDA内核的块和网格大小
dim3 blockSize(1, 256);
const int iterations = 10000;
size_t imgSize = width * height * sizeof(uchar3);
size_t graySize = width * height * sizeof(unsigned char);
// 在GPU上分配内存
cudaMalloc(&d_img, imgSize);
cudaMalloc(&d_gray, graySize);
// 将图像数据从CPU复制到GPU
cudaMemcpy(d_img, img.data, imgSize, cudaMemcpyHostToDevice);
// 取上整除以确保覆盖所有像素
dim3 gridSize((width + blockSize.x - 1) / blockSize.x, (height + blockSize.y - 1) / blockSize.y);
// 启动CUDA内核进行RGB到灰度的转换
rgbToGray<<<gridSize, blockSize>>>(reinterpret_cast<uchar3*>(d_img), d_gray, width, height);
// 检查内核启动是否成功
cudaGetLastError();
// 同步
cudaDeviceSynchronize();
// 测试时间
cudaEvent_t start, stop;
float total_time = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cout << "正在测试GPU性能..." << endl;
for (int i = 0; i < iterations; i++) {
cudaEventRecord(start);
rgbToGray<<<gridSize, blockSize>>>(reinterpret_cast<uchar3*>(d_img), d_gray, width, height);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float single_time = 0.0f;
cudaEventElapsedTime(&single_time, start, stop);
total_time += single_time;
}
float avg_time = total_time / iterations;
cout << "平均每次转换时间: " << avg_time << " ms" << endl;
// 将结果从GPU复制回CPU
cv::Mat grayImg(height, width, CV_8UC1);
cudaMemcpy(grayImg.data, d_gray, graySize, cudaMemcpyDeviceToHost);
// 保存灰度图像
cv::imwrite("output.png", grayImg);
// 释放GPU内存
cudaFree(d_img);
cudaFree(d_gray);
return 0;
}

2.3 三维坐标计算#

假设有一个 256 × 256 × 128 的CT图像(宽×高×深度),每个体素是一个16位整数。我们要用CUDA把它变成浮点数,并除以最大灰度值,得到归一化的3D数组。

图示#

三维坐标计算演示
三维坐标计算演示

内存布局#

我们同样把三维数据展平成一维数据。

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < dimX && y < dimY && z < dimZ) {
// 因为内存布局是行优先,而x是变化最快的,因此这种访问方式能够让线程访问连续的x,能够合并内存访问,性能较好。
int global_idx = z * dimY * dimX + y * dimX + x;
}
三维块的硬件限制

CUDA 规定每个 Block 的总线程数不能超过 1024。 三维块尺寸乘积必须 ≤ 1024,例如 8×8×8 = 512 合法,16×16×16 = 4096 非法。

核心的Kernel编写如下:

__global__ void normalizeVolume(const unsigned short* d_in,
float* d_out,
int dimX, int dimY, int dimZ,
float maxVal) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < dimX && y < dimY && z < dimZ) {
int idx = z * dimY * dimX + y * dimX + x;
d_out[idx] = (float)d_in[idx] / maxVal;
}
}

三、启动Kernel#

算出了索引,还得告诉GPU启动多少个线程,怎么分组?

这就是<<<gridDim, blockDim>>>的作用。

搞懂内置变量和类型#

CUDA内核函数__global__中可以直接使用以下内置变量,无需声明。

之前我们看到的threadIdx,blockIdx,blockDim,gridDim实际上它们的类型都是dim3

dim3是CUDA内置的一个结构体,包含xyz三个无符号整数字段。

无论多少维度的公式,核心公式都是:

全局线程ID = blockIdx维度 * blockDim维度 + threadIdx维度

当数据是一维时,只需要算x即可;当二维或者三维,则需要对其进行拓展。

因此我们在启动Kernel的时候,需要对其进行配置。

/* 一维配置 */
int threadsPerBlock = 256;
// 向上取整
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
kernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, N);
/* 二维配置 */
dim3 blockSize(16, 16); // 16×16 = 256 线程
dim3 gridSize( (width+15)/16, (height+15)/16 );
kernel<<<gridSize, blockSize>>>(d_image, width, height);
/* 三维配置 */
dim3 blockSize(8, 8, 4); // 8×8×4 = 256 线程 ≤ 1024
dim3 gridSize( (dimX+7)/8, (dimY+7)/8, (dimZ+3)/4 );
kernel<<<gridSize, blockSize>>>(d_volume, dimX, dimY, dimZ);
关于dim3的易错点
  • 没有赋值的维度默认为 1。例如 dim3 block(32, 32) → block.z = 1。
  • 内核中访问 threadIdx.z 是安全的,但如果启动时没给 z 值(或给了1),它始终为 0。
  • 用整数直接传给 <<<>>> 时,相当于只设置了 .x 分量。例如 kernel<<<10, 256>>> 等价于 gridDim.x=10, blockDim.x=256,其他维为1。

参数解释#

  • threadsPerBlock:每个 Block 的线程数,通常取 32 的整数倍(如 128、256、512),避免 Warp 资源浪费。
  • blocksPerGrid向上取整确保覆盖全部数据。公式等价于 ceil(N / threadsPerBlock)

为什么二维配置常用 16×16 或 32×8?

  • 合并访问:x 维度的线程连续访问内存,性能最优。
  • Warp 对齐:32 个线程为一个 Warp,块尺寸最好是 Warp 大小的倍数。
  • 共享内存:二维块切出的瓦片(tile)更规整,便于利用共享内存。

思考题:16×16 和 32×8 的块,哪个更适合处理 1024×1024 的图像?
提示:考虑合并访问和 Warp 利用率(32 个线程为一组)。32×8 的 x 维有 32 个线程,正好一个 Warp,合并访问效率最高。

四、一维、二维、三维的选择策略#

数据形状推荐网格/块维度理由
一维数组一维索引计算最简单,开销最小
二维图像/矩阵二维块 + 二维网格直观,便于二维 tile 划分,合并访问最优
三维体数据三维块 + 三维网格逻辑清晰,利于 3D 局部性
多层二维数据二维网格 + 循环用二维网格处理每层,代码复用性高

黄金法则让 x 维度的线程连续访问连续的内存地址。
因为一个 Warp 内的 32 个线程是沿 x 维度连续的,合并访问是免费的性能提升。

小结#

  • 一维索引blockIdx.x * blockDim.x + threadIdx.x
  • 二维索引:分开计算 rowcol,注意行优先内存排布
  • 三维索引:依次展开 zyx
  • 配置原则:块尺寸取 32 的倍数,网格数量向上取整覆盖全部数据
  • 选择依据:数据形状 + 内存合并访问

记住:CUDA 的坐标计算本质上就是“基地址 + 偏移量”的变体。
理解了这一点,任何维度的索引你都能自己推导出来。

参考文献#

支持与分享

如果这篇文章对你有帮助,欢迎分享给更多人或赞助支持!

赞助
CUDA学习之路[4]——CUDA全局坐标计算
https://dlog.com.cn/posts/cuda04/cuda全局坐标计算/
作者
杜子源
发布于
2026-03-28
许可协议
CC BY-NC-SA 4.0
最后更新于 2026-03-28,距今已过 56 天

部分内容可能已过时

Profile Image of the Author
杜子源
都是风景,幸会
公告
请狠狠地打赏我,打赏一次,爆更一篇!!
音乐
封面

音乐

暂未播放

0:00 0:00
暂无歌词
分类
标签
站点统计
文章
26
分类
8
标签
11
总字数
66,827
运行时长
0
最后活动
0 天前

目录