【CUDA手册004】一个典型算子的 CUDA 化完整流程

在本篇中，我们将以医学图像处理中最基础的“二值化阈值算子（Thresholding）”为例，演示如何将一个 C++ 算子完整地迁移到 CUDA。

1. 第一步：建立 CPU Baseline（基准）

在写任何 CUDA 代码前，必须先有一个性能表现稳定、逻辑正确的 CPU 版本。这不仅是为了比对结果，更是为了后续做性能分析。

// CPU 版本：简单的阈值处理voidThresholdCPU(constfloat*input,float*output,intwidth,intheight,floatthreshold){intsize=width*height;for(inti=0;i<size;++i){output[i]=(input[i]>threshold)?1.0f:0.0f;}}

2. 第二步：编写 CUDA Kernel（核心计算）

按照第二篇讲到的映射逻辑，我们将for循环拆解到每一个线程中。

// CUDA Kernel__global__voidThresholdKernel(constfloat*input,float*output,intsize,floatthreshold){intidx=blockIdx.x*blockDim.x+threadIdx.x;if(idx<size){// 这里的逻辑与 CPU 版本完全一致output[idx]=(input[idx]>threshold)?1.0f:0.0f;}}

3. 第三步：工程化封装（Host 侧逻辑）

不要在业务代码里到处写cudaMalloc。我们需要一个清晰的封装函数，处理内存申请、拷贝、启动和释放。

// 工程化封装：隐藏 CUDA 细节voidThresholdGPU(constfloat*h_input,float*h_output,intwidth,intheight,floatthreshold){intsize=width*height;size_t byteSize=size*sizeof(float);float*d_input,*d_output;// 1. 资源分配cudaMalloc(&d_input,byteSize);cudaMalloc(&d_output,byteSize);// 2. 数据上传cudaMemcpy(d_input,h_input,byteSize,cudaMemcpyHostToDevice);// 3. 执行配置intthreadsPerBlock=256;intblocksPerGrid=(size+threadsPerBlock-1)/threadsPerBlock;// 4. 启动算子ThresholdKernel<<<blocksPerGrid,threadsPerBlock>>>(d_input,d_output,size,threshold);// 5. 数据回传cudaMemcpy(h_output,d_output,byteSize,cudaMemcpyDeviceToHost);// 6. 释放资源cudaFree(d_input);cudaFree(d_output);}

4. 第四步：性能对比与落地

医学图像通常很大（例如 4K 的 X 光片或大尺度的 CT 序列）。我们必须通过高精度计时器来验证“落地”效果。

实验数据参考（以 4096 x 4096 图像为例）：

• CPU (i7-12700K)：约 12ms
• GPU (RTX 3060, 含拷贝时间)：约 3ms
• GPU (RTX 3060, 纯计算时间)：约 0.1ms

为什么含拷贝时间后提升没那么大？

对于阈值处理这种**计算强度极低（Arithmetic Intensity low）**的算子，PCIe 总线的带宽往往是瓶颈。如果你的算法流程里只有这一个 GPU 算子，那么把数据搬来搬去可能并不划算。真正的 CUDA 高手会把多个算子串联在 GPU 上，只在开头和结尾做数据拷贝。

5. 数据类型与归一化

在真实的医学场景中，数据往往不是float，而是uint16（CT 的 HU 值）或uint8。

• 对齐问题：如果处理uint16数据，要注意内存对齐。
• 动态范围：医学图像通常有很高的动态范围（如 0-4095），在做归一化（Normalization）时，需要先在 GPU 上找到最大最小值。

6. 产出能力：第一个 CUDA 算子类

建议在工程中开始建立如下结构：

Project/ ├── include/ │ └── ImageOps.hpp (声明 CPU/GPU 接口) ├── src/ │ ├── Threshold.cpp (CPU 实现) │ └── Threshold.cu (GPU 实现与 Wrapper) └── test/ └── main.cpp (比对结果：assert(gpu_res == cpu_res))