CUDA Programming Model
内容提要
本文介绍CUDA编程模型：异构计算中GPU处理计算密集任务，CPU负责其余部分。核心概念包括Host与Device、执行空间说明符（__global__、__device__、__host__）、线程层次（Thread、Block、Grid）及内置变量。还讲解NVCC两段式编译流程（虚拟架构与真实SM架构）及用NVProf分析GPU性能。
延伸解读
线程层次设计的实际意义
文章指出，Block 的引入并非多余，而是为了支持线程协作和扩展性。同一 Block 内的线程可协作，且 Block 可独立调度，避免全局同步。这种设计允许在增加 SM 时，旧代码也能高效加速，同时限制协作范围，减少内存带宽压力。理解这一点有助于编写可扩展的 CUDA 程序。
执行空间说明符的使用限制
文章强调，__global__ 函数必须返回 void，不能是类成员，且调用需指定执行配置。__device__ 和 __host__ 可同时使用，为主机和设备编译函数。但 __global__ 不能与 __host__ 或 __device__ 共用。这些规则影响函数定义和调用方式，编写时需严格遵守。
NVCC 两段式编译的兼容性考量
为兼容不同 GPU 架构，NVCC 采用虚拟架构和真实 SM 架构两段式编译。虚拟架构的计算能力不能大于真实 SM 架构。编译时需指定 compute_XX 和 sm_XX，如 nvcc x.cu --gpu-architecture=compute_50 --gpu-code=sm_50,sm_52。这确保生成的代码能在目标 GPU 上运行，并可能利用新架构特性。
NVProf 性能分析工具的作用
NVProf 是 CUDA Toolkit 自带工具，可生成 GPU 时间线，帮助识别 GPU 闲置或利用不足。通过 --print-gpu-trace 可查看内核执行和内存复制的详细信息，如内核参数、共享内存使用和内存吞吐量。结合 Nvvp 或 Nsight 可视化，能更准确定位性能瓶颈，指导优化。
Q&A
CUDA中的Host和Device分别指什么？
Host指CPU和主机内存（host memory），Device指GPU和显存（device memory）。
__global__、__device__和__host__这三个执行空间说明符有什么区别？
__global__声明内核函数，在设备上执行，可从主机调用，必须返回void；__device__声明的函数在设备上执行，只能从设备调用；__host__声明的函数在主机上执行，只能从主机调用。__global__不能与__host__或__device__同时使用，但__device__和__host__可以一起使用。
CUDA的线程层次结构是怎样的？
CUDA线程层次分为Thread、Thread Block和Thread Grid。Thread是顺序执行单元，所有线程并行执行相同的核函数；Block是一组线程，执行在一个SM上，同一Block中的线程可以协作；Grid是Block的集合，一个Grid中的Block可以在多个SM中执行。
CUDA的内置变量threadIdx、blockIdx、blockDim和gridDim分别代表什么？
threadIdx是当前线程在block中的索引；blockIdx是当前block在grid中的索引；blockDim表示一个block中包含的线程数；gridDim表示一个grid中包含的block数。它们都有x、y、z三个维度。
NVCC编译CUDA程序时，虚拟架构和真实SM架构是什么？编译时要注意什么？
NVCC采用两段式编译：虚拟计算架构（compute_xx）和真实SM架构（sm_xx）。编译时要根据GPU架构选择合适的计算能力，且虚拟架构的计算能力不能大于真实SM架构，即arch <= code。例如：nvcc x.cu --gpu-architecture=compute_50 --gpu-code=sm_50,sm_52。
如何使用NVProf分析CUDA程序的性能？
NVProf是NVIDIA提供的GPU性能分析工具，可以生成GPU时间线。常用命令：nvprof a.exe 运行程序并输出分析；nvprof -o out.nvvp a.exe 输出到文件，可结合Nvvp或Nsight可视化；nvprof --print-gpu-trace a.exe 显示GPU活动时间线；nvprof --print-api-trace a.exe 显示主机端CUDA API调用时间线。
