4、CUDA编程模型（下）：核函数、内存管理与同步机制
好，咱们接着聊。上一章我们把CUDA的线程层次结构和网格、块的概念理清了。这一章，我带你深入实战——写核函数、管好内存、搞定同步。这些是CUDA编程的硬功夫，也是我当年踩坑最多的地方。
4.1 核函数编写：从入门到实战
核函数，说白了就是在GPU上跑的函数。写法上有个硬性规定：前面必须加 __global__ 修饰符，返回值必须是 void。我刚开始写的时候，老忘记加这个修饰符，结果编译报错半天找不到原因。
__global__ void vecAdd(float* A, float* B, float* C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
C[i] = A[i] + B[i];
}
}你看这个经典的向量加法。每个线程只处理一个元素。这里有个细节：if (i < N) 这个边界检查，我建议你永远别省。为什么？因为你的线程总数（网格大小×块大小）不一定能整除数据长度，多出来的线程必须跳过，否则会访问非法内存。
核心要点：核函数里不能调用标准C库函数（比如printf在调试模式下可以，但正式代码别用），也不能使用静态变量。每个线程都是独立的执行流。
4.2 内存管理：cudaMalloc 与 cudaMemcpy
GPU和CPU是两套独立的内存系统。你在CPU上malloc出来的指针，GPU核函数是访问不到的。必须用 cudaMalloc 在显存上分配空间。
我见过太多新手直接传CPU指针进核函数，然后程序崩溃。嗯，这其实是个经典错误。
float *d_A, *d_B, *d_C;
int size = N * sizeof(float);
// 在GPU上分配内存
cudaMalloc((void**)&d_A, size);
cudaMalloc((void**)&d_B, size);
cudaMalloc((void**)&d_C, size);
// 把数据从CPU拷贝到GPU
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// 调用核函数
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
// 把结果拷回CPU
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// 别忘了释放显存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);个人经验：我习惯在每次cudaMalloc和cudaMemcpy之后加错误检查。用 cudaGetLastError() 或者 cudaError_t 判断返回值。别嫌麻烦，这能帮你省下大量调试时间。
这里有个表格，帮你快速记住三种传输方向：
| 枚举值 | 含义 | 使用场景 |
|---|---|---|
| cudaMemcpyHostToDevice | CPU → GPU | 输入数据上传 |
| cudaMemcpyDeviceToHost | GPU → CPU | 结果回传 |
| cudaMemcpyDeviceToDevice | GPU → GPU | 显存内部拷贝 |
4.3 同步机制：__syncthreads()
同一个块内的线程，执行进度可能不一样。有的线程跑得快，已经算完了；有的线程还在等数据。这时候如果你让快的线程去读慢的线程正在写的数据，结果就是错的。
__syncthreads() 就是用来解决这个问题的。它像一个栅栏，所有线程必须都到达这个点，才能继续往下走。
__global__ void sharedMemoryExample(float* data) {
__shared__ float sdata[256];
int tid = threadIdx.x;
sdata[tid] = data[tid];
__syncthreads(); // 等所有线程都写完共享内存
// 现在可以安全地读其他线程写的数据了
if (tid < 128) {
data[tid] = sdata[tid] + sdata[tid + 128];
}
}我曾经踩过的坑：在条件分支里用__syncthreads()。比如if(tid < 32) { __syncthreads(); }。这会导致死锁！因为只有部分线程执行了同步，其他线程永远等不到。记住：__syncthreads()必须被块内所有线程执行到。
4.4 流与事件：让GPU忙起来
默认情况下，CUDA操作是串行的。你发一个核函数，等它执行完，再发下一个。但GPU有那么多计算单元，完全可以同时干好几件事。
流（Stream）就是用来实现并发的。你可以把不同的操作放到不同的流里，让它们重叠执行。
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 两个流可以并行执行
kernel1<<<grid, block, 0, stream1>>>(d_A);
kernel2<<<grid, block, 0, stream2>>>(d_B);
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
cudaStreamDestroy(stream1);
cudaStreamDestroy(stream2);事件（Event）则用来做时间测量和同步控制。你可以记录某个操作什么时候完成，然后让另一个操作等它。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
kernel<<<grid, block>>>(d_data);
cudaEventRecord(stop);
cudaEventSynchronize(stop); // 等待事件完成
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("核函数耗时: %f ms\n", milliseconds);
cudaEventDestroy(start);
cudaEventDestroy(stop);我的建议：性能调优时，先用事件测量每个阶段的耗时。找到瓶颈在哪，再针对性优化。别凭感觉猜，数据说话。
4.5 本章知识体系
下面这张图，帮你把本章的核心逻辑串起来：
你看，这四个模块是环环相扣的。核函数是执行体，内存管理是数据基础，同步保证正确性，流与事件提升效率。少了哪一个，你的CUDA程序都跑不顺畅。