1
int i = blockIdx.x * blockDim.x + thredIdx.x

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

量 含义 维度范围
threadIdx.x/y/z 线程在块内的局部索引 日 blockDim-1
blockIdx.x/y/z 线程块在网格内的索引 日 gridDim-1
blockDim.x/y/z 每个块每维的线程数 由启动参数 决定 $<<\ldots,\quad{\mathrm{t h r e a d s}}>>>$
gridDim.x/y/z 网格每维的块数 由启动参数 <<

所谓全局索引，就是跳过前面所有块的线程，再加上当前块内的偏移。

全局坐标计算方式

在内存中，数据始终按照一维排布，因此全局坐标计算就是将多维逻辑索引映射为一维物理地址

一维坐标计算

1
int global_id = blockIdx.x * blockDim.x + threadIdx.x;
  • blockIdx * blockDim.x 计算之前的块有多少线程（基地址），用来定位到当前块
  • ‘threadIdx.x’ 加上当前线程块内的偏移

边界保护

1
2
3
if (global_id < N){
    // 安全处理
}
  • 为什么要这个判断？
    • 网格覆盖线程总数是块大小整数倍，可能超过数据总量N，不加判断会越界

二维坐标计算

二维数据（比如图像）需要两个全局索引: ‘row’ 和 ‘col’

1
2
3
4
5
6
7
8
9
10
// 列方向，当前块之前的数量 * 每块宽度 + 当前块内的线程偏移
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图像转灰度图

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
// 输入图像在显存中是按 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，确保覆盖所有像素

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
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);

    // ... 后续同步与结果保存 ...
}

完整代码

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
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
#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;
}

三维坐标计算

示例：图像转化归一化

将一个 256 x 256 x 128 CT图像变成浮点数，除以最大灰度值，得到归一化3D数组

内存布局

将三维数据展平为一维数据

1
2
3
4
5
6
7
8
9
10
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
    核心kernel编写如下：
    1
    2
    3
    4
    5
    6
    7
    8
    9
    10
    11
    12
    13
    __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__ 中可以直接使用一下内置变量，无需声明.
比如之前的treadIdx, blockIdx, blockDim, gridDim 实际上的类型都是 dim3，dim3是CUDA内置的一个结构体，包含xyz三个无符号整数字段。
无论维度为几，核心公式都是：

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

如果数据是一维，只需要计算x；如果是二维或三维，则需要对其进行扩展。
因此启动Kerrnel时，需要对其进行配置

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
/* 一维配置 */
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);
  • 易错点
    • 没有赋值的维度默认为1
    • 内核中访问threadIdx.z是安全的，但如果启动时没给z值，它始终为0
    • 用整数直接传给 <<<<>>>> 时，相当于只设置.x分量。例如 kernel<<<10,256>>>等价于 gridDim.x=10,blockDim.x=256，其他维为1

参数解释

  • treadsPerBlock:每个block的线程数，一般取32倍数防止warp资源浪费
  • blocksPerGrid：向上取整确保覆盖全部数据，公式等价与 $\text{ceil}(N / threadsPerBlock)$
    为什么二维配置一般使用16 x 166或者 32 x 8
    • 合并访问：x 维度的线程连续访问内存，性能最优
    • warp对齐：32个线程为一个Warp，块尺寸最好时Warp大小倍数
    • 共享内存：二维块切除的tile更规整，便于利用共享内存

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

数据形状 推荐网格/块维度 理由
一维数组 一维 索引计算最简单，开销最小
二维图像/矩阵 二维块+二维网格 直观，便于二维tile划分，合并访问最优
三维体数据 三维块+三维网格 逻辑清晰，利于3D局部性
多层二维数据 二维网格+循环 用二维网格处理每层，代码复用性高
  • 黄金法则：让 x 维度的线程连续访问连续的内存地址， 一个Warp内的32个线程时沿x维度连续的，合并访问时免费的性能提升