Writing Your First CUDA Program
内容提要
本文介绍第一个CUDA程序：用__global__定义核函数，用<<<>>>配置线程，用cudaDeviceSynchronize()同步；演示nvcc编译、编写Makefile及多文件编译；讲解threadIdx、blockIdx等线程索引；最后用nvprof分析性能。
延伸解读
异步执行与同步的必要性
CUDA 内核启动是异步的，CPU 在启动内核后立即继续执行，不会等待 GPU 完成。如果主机程序随后退出，内核可能还未执行，导致输出丢失。因此，在启动内核后调用 cudaDeviceSynchronize() 是确保 GPU 任务完成并获取结果的常见做法。这一点在第一个示例中就已体现，是理解 CUDA 执行模型的基础。
从单文件到多文件编译的过渡
文章展示了从简单 nvcc 编译到使用 Makefile 管理编译的过程，并进一步扩展到多文件协同编译。多文件编译时，需要将包含核函数的文件用 --device-c 选项编译为对象文件，再与其他源文件链接。这反映了实际项目中代码组织与构建的常见需求，也为后续集成第三方库（如 cuBLAS）提供了基础。
线程索引：定位并行执行单元
在 CUDA 中，每个线程通过 threadIdx 和 blockIdx 等内置变量唯一标识。threadIdx.x 表示线程在块内的序号，blockIdx.x 表示线程块在网格中的序号，而 blockDim.x 和 gridDim.x 分别描述块和网格的维度。示例中启动 <<<33, 5>>> 的配置，输出显示了不同块和线程的编号，帮助直观理解线程的组织方式。
性能分析工具 nvprof 的初步使用
文章最后使用 nvprof 对程序进行性能分析，输出中区分了 GPU 上内核运行时间（Profiling result）和 CPU 上 API 调用时间（API calls）。这有助于开发者了解程序在 GPU 和 CPU 上的耗时分布，为优化提供依据。nvprof 是 CUDA 工具包中的性能分析工具，适合入门阶段观察内核执行情况。
Q&A
如何编写第一个CUDA程序？
创建一个.cu文件，用__global__修饰核函数，在main中调用时使用<<<1,1>>>配置线程，并调用cudaDeviceSynchronize()等待GPU完成。例如：__global__ void hello_from_gpu() { printf("Hello World from the GPU!\n"); } int main() { hello_from_gpu<<<1,1>>>(); cudaDeviceSynchronize(); return 0; }
为什么CUDA核函数调用后需要cudaDeviceSynchronize()？
因为内核启动是异步的，CPU在启动GPU进程后立即继续执行，如果不等待，应用程序可能直接退出，导致内核输出无法显示。cudaDeviceSynchronize()会阻塞CPU直到GPU完成，确保输出可见。
如何用nvcc编译CUDA程序？
使用命令：/usr/local/cuda/bin/nvcc hello_cuda.cu -o hello_cuda -run，其中-run选项表示编译后立即运行。
如何编写Makefile来编译CUDA程序？
可以编写Makefile，指定源文件、目标文件和编译器。例如：TEST_SOURCE = hello_cuda.cu TARGETBIN := ./hello_cuda CC = /usr/local/cuda/bin/nvcc $(TARGETBIN):$(TEST_SOURCE) $(CC) $(TEST_SOURCE) -o $(TARGETBIN) .PHONY:clean clean: -rm -rf $(TARGETBIN) -rm -rf *.o 然后运行make即可。
CUDA中如何获取线程索引？
使用threadIdx.x获取线程在块内的x方向索引，blockIdx.x获取线程块在grid中的x方向索引，blockDim.x获取块内x方向线程数，gridDim.x获取grid中x方向块数。例如：const int bid = blockIdx.x; const int tid = threadIdx.x;
如何用nvprof分析CUDA程序性能？
使用命令：/usr/local/cuda/bin/nvprof ./Index_of_thread，它会输出Profiling result（GPU上kernel运行时间）和API calls（CPU上API调用时间）。
