2、CUDA基础回顾：GPU架构、内存层次、线程模型、CUDA编程模型
在深入排查CUDA与TensorRT推理异常之前，必须先建立对GPU底层执行模型的清晰认知。绝大多数推理异常——无论是非法内存访问、未初始化值、还是性能骤降——其根源都能追溯到对GPU架构或编程模型的误解。本节将系统梳理四个核心维度：GPU硬件架构、内存层次、线程组织模型，以及CUDA编程的基本范式。
2.1 GPU硬件架构：从SM到CUDA Core
NVIDIA GPU的物理计算单元以流式多处理器（Streaming Multiprocessor, SM）为基本单位。每个SM内部包含若干CUDA Core（FP32单元）、Tensor Core（混合精度矩阵运算单元）、特殊函数单元（SFU）、以及共享内存和寄存器堆。以Ampere架构的GA102为例，每个SM包含128个CUDA Core和4个Tensor Core。
理解SM的调度机制对排查性能问题至关重要：SM以warp（32线程）为基本调度单位，硬件在每个时钟周期内选择一个就绪的warp发射指令。当某个warp因访存或同步等待而阻塞时，SM会立即切换到其他就绪warp，这种延迟隐藏机制是GPU高吞吐量的核心保障。若kernel的并行度不足以填满SM的warp槽位，或每个线程的寄存器使用过多导致占用率下降，GPU就无法有效隐藏访存延迟，表现为推理延迟异常升高。
2.2 内存层次：延迟与可见性的权衡
CUDA的内存模型是分层级的，每一层在访问速度和可见范围上都有严格区别。理解这些区别是排查“数据不对”类问题的前提。
| 内存类型 | 作用域 | 生命周期 | 典型用途 |
|---|---|---|---|
| 寄存器 | 单个线程 | 线程执行期间 | 局部变量、中间计算 |
| 共享内存 | 单个Block | Block执行期间 | Block内数据复用、协作 |
| 全局内存 | 所有线程 | 由cudaMalloc分配，直到cudaFree | 主数据存储、跨Block通信 |
| 常量内存 | 所有线程（只读） | 程序运行期间 | 只读参数、权重系数 |
| 纹理/表面内存 | 所有线程（特殊访问模式） | 由CUDA数组管理 | 图像处理、插值访问 |
在TensorRT推理场景中，最常见的错误之一是在kernel中直接访问全局内存而未利用共享内存做数据复用，导致访存带宽成为瓶颈。另一个典型问题是未对齐访问：当线程访问的全局地址不满足128字节对齐时，硬件会发起多次内存事务，造成有效带宽骤降。排查时可通过ncu --metrics dram__bytes_read.sum观察实际吞吐与理论峰值的差距。
2.3 线程模型：Grid、Block、Warp与Thread
CUDA采用三层线程组织：Grid（网格）→ Block（线程块）→ Thread（线程）。硬件层面，Block会被进一步划分为32线程一组的Warp执行。这个模型直接决定了kernel的并行度上限和资源分配方式。
// 典型的kernel启动配置
dim3 blockDim(256); // 每个Block 256个线程
dim3 gridDim((N + 255) / 256); // 向上取整覆盖N个元素
myKernel<<<gridDim, blockDim, sharedMemSize, stream>>>(d_data, N);排查异常时，需要关注三个关键参数：
- Block大小：通常取32的倍数（如128/256/512）。过小则无法隐藏延迟，过大则可能超出SM的线程上限（如1024）导致启动失败。
- Grid大小：必须覆盖全部数据。若gridDim计算有误（如整数除法截断），会导致部分数据未被处理，表现为输出中随机位置的错误值。
- 共享内存大小：动态共享内存需在启动时指定，若超过SM剩余容量，kernel会启动失败并返回
cudaErrorInvalidValue。
2.4 CUDA编程模型：Host与Device的协作
CUDA程序分为Host端（CPU）和Device端（GPU）两部分。Host负责数据准备、kernel启动和结果回收；Device负责并行计算。两者通过PCIe总线通信，而数据传输往往是推理性能的隐藏瓶颈。
标准流程如下：
- 分配Host内存并填充输入数据
- 使用
cudaMalloc分配Device全局内存 - 使用
cudaMemcpy将输入从Host拷贝到Device - 配置kernel启动参数并执行
- 使用
cudaMemcpy将结果拷回Host - 释放所有分配的内存
在TensorRT中，这一流程被引擎封装，但底层仍然遵循相同模型。排查推理异常时，需要特别关注异步操作的正确性：
cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_input, h_input, bytes, cudaMemcpyHostToDevice, stream);
myKernel<<<grid, block, 0, stream>>>(d_input, d_output, N);
cudaMemcpyAsync(h_output, d_output, bytes, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream); // 必须同步后才能安全读取h_outputcudaMemcpyAsync后未同步就读取Host端结果，可能读到未更新的旧数据。使用cudaDeviceSynchronize()或cudaStreamSynchronize()确保所有异步操作完成。此外，cudaMemcpy返回cudaErrorInvalidValue时，优先检查拷贝字节数是否与分配大小匹配。
2.5 案例复盘：一次典型的越界访问
某推理服务在运行TensorRT FP16引擎时，输出张量每隔固定间隔出现NaN。初步排查排除了模型权重问题，最终定位到自定义plugin kernel的索引计算错误。
原始代码如下：
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 错误：未检查idx是否超出输出范围
if (idx < outputSize) {
float val = input[idx * 2] + input[idx * 2 + 1];
output[idx] = val;
}问题在于input数组长度为outputSize * 2，但kernel启动时gridDim计算为(outputSize + 255) / 256，当outputSize不是256的倍数时，最后一个Block的部分线程会访问input[outputSize * 2]及之后的越界地址。虽然条件判断idx < outputSize保护了output的写入，但input的读取发生在判断之前，导致越界读取。
修复方案：
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < outputSize) {
int inIdx = idx * 2;
// 确保inIdx+1不越界
if (inIdx + 1 < inputSize) {
float val = input[inIdx] + input[inIdx + 1];
output[idx] = val;
}
}更稳妥的做法是使用cuda-memcheck或compute-sanitizer工具，在开发阶段即可捕获越界访问，避免线上故障。