CUDA学习之路[4]——CUDA全局坐标计算
目录
引言
在上一篇中,我们理清了 GPU 的硬件架构与 CUDA 的线程模型:Grid → Block → Thread。
但很多同学一到写代码就卡住,最先卡住的就是算不对全局索引。
int i = blockIdx.x * blockDim.x + threadIdx.x;- 彻底掌握一维、二维、三维全局索引的计算公式
- 理解
gridDim、blockDim与启动配置之间的对应关系 - 学会根据数据形状选择合适的网格/块维度
思考题:假如你有一个 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 二维坐标计算
二维数据(如图像)需要两个全局索引:row 和 col。
公式
// 列方向,当前块之前的数量 * 每块宽度 + 当前块内的线程偏移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来将彩色图片转为灰度图。
sudo apt updatesudo 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 线程 ≤ 1024dim3 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 - 二维索引:分开计算
row和col,注意行优先内存排布 - 三维索引:依次展开
z、y、x - 配置原则:块尺寸取 32 的倍数,网格数量向上取整覆盖全部数据
- 选择依据:数据形状 + 内存合并访问
记住:CUDA 的坐标计算本质上就是“基地址 + 偏移量”的变体。
理解了这一点,任何维度的索引你都能自己推导出来。
参考文献
支持与分享
如果这篇文章对你有帮助,欢迎分享给更多人或赞助支持!
部分内容可能已过时