4、核函数编写基础：修饰符、参数传递与执行控制
好，咱们正式开始写核函数了。
说实话，很多初学者一上来就盯着 __global__ 看，觉得这就是CUDA的全部。其实不然。修饰符只是门面，真正决定性能的，是你怎么组织代码、怎么传参数、怎么控制执行流。今天这一章，我把这些基础但关键的点，掰开揉碎了讲清楚。
4.1 三个修饰符：各司其职
CUDA 里就三个修饰符：__global__、__device__、__host__。别小看它们，用错了轻则编译报错，重则性能翻车。
__global__：从主机调用，在设备上执行。返回值必须是void。这是核函数的入口。__device__：从设备调用，在设备上执行。可以返回非 void 类型。说白了，就是 GPU 上的普通函数。__host__：从主机调用，在主机上执行。默认就是它，所以一般不用写。
关键点：__global__ 不能和 __device__ 或 __host__ 同时使用。但 __device__ 和 __host__ 可以同时出现，表示这个函数既能在 CPU 跑，也能在 GPU 跑。
我个人习惯是，把频繁调用的计算逻辑写成 __device__ 函数。比如矩阵乘法里的内积计算，写成 __device__ float dot(...)，然后在核函数里反复调用。这样代码清晰，也方便单元测试。
小技巧：如果你不确定某个函数该放哪，先写成 __host__ __device__。编译时 CUDA 会帮你生成两份代码。代价是代码体积变大，但调试期很实用。
4.2 核函数参数传递：值传递 vs 指针传递
核函数的参数传递，其实和 C 语言一样。但有个坑：参数是传值的。你传一个结构体进去，CUDA 会把它复制到每个线程的寄存器或本地内存里。结构体太大？那寄存器就爆了，性能直接崩。
我曾经在一个项目里，传了一个 256 字节的结构体给核函数。结果每个线程都在拼命挤寄存器， occupancy 掉到 30% 以下。后来改成传指针，性能翻了 3 倍。
所以我的建议是：
- 传 基本类型（int, float 等）：直接传值，没问题。
- 传 数组或大结构体：传指针，让线程通过指针去全局内存读。
- 传 小结构体（<= 16 字节）：可以传值，但要注意对齐。
// 推荐：传指针
__global__ void vec_add(float *a, float *b, float *c, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) c[idx] = a[idx] + b[idx];
}
// 不推荐：传大结构体
struct BigData { float data[64]; };
__global__ void bad_kernel(BigData d) { ... } // 每个线程都复制一份，寄存器爆炸注意：核函数参数最多 512 字节（计算能力 2.x 以上是 4KB）。但别真用到上限，参数越多，启动开销越大。我一般控制在 8 个参数以内。
4.3 异步执行：别让 GPU 等你
CUDA 默认是异步的。你调用 kernel<<<grid, block>>>() 之后，控制权立刻返回给 CPU。CPU 可以继续做别的事，比如准备下一批数据。
但有个陷阱：异步只对同一个流（stream）有效。如果你不指定流，默认用默认流（stream 0）。默认流在某些情况下会阻塞其他操作。
我刚开始做训练时，就踩过这个坑。CPU 在等 GPU 完成，GPU 在等 CPU 发数据，两边互相等，利用率不到 50%。后来用 cudaStreamCreate() 创建多个流，把数据拷贝和计算重叠起来，吞吐量直接翻倍。
异步执行的核心模式：
- CPU 启动核函数（异步）
- CPU 立即返回，准备下一批数据
- GPU 在后台执行计算
- CPU 在需要结果时，调用
cudaDeviceSynchronize()等待
cudaMemcpyAsync(d_x, h_x, size, cudaMemcpyHostToDevice, stream1);
kernel<<<grid, block, 0, stream1>>>(d_x, d_y);
cudaMemcpyAsync(h_y, d_y, size, cudaMemcpyDeviceToHost, stream1);
// CPU 可以在这里做其他事
cudaDeviceSynchronize(); // 等所有流完成核心原则：能用异步就别用同步。把 cudaMemcpy 换成 cudaMemcpyAsync，把核函数绑定到非默认流上。这是性能优化的第一步。
4.4 同步：什么时候必须等？
同步不是坏事，但滥用同步就是灾难。常见的同步点：
cudaDeviceSynchronize()：等所有流上的所有操作完成。最重，慎用。cudaStreamSynchronize(stream)：等某个流完成。轻量一些。__syncthreads()：块内线程同步。用于共享内存读写协调。
我见过有人每个核函数后面都加 cudaDeviceSynchronize()，说是「为了保险」。结果 GPU 利用率一直在 20% 以下。其实只要在真正需要结果的地方同步一次就够了。
举个例子：训练一个 batch，你可以在前向、反向、权重更新之间同步一次，而不是每层都同步。
避坑指南：我曾经在调试时，发现结果不对，就到处加同步。结果问题没找到，性能先崩了。后来用 cuda-memcheck 才发现是越界访问。同步不能解决逻辑错误，它只解决时序问题。
4.5 知识体系总览
下面这张图，把本章的核心逻辑串起来了。你可以把它当作核函数编写的「路线图」。
这张图里，三个分支最终都指向「实践建议」。说白了，修饰符、参数传递、执行控制，最终都是为了写出高效、稳定的核函数。没有哪一个是孤立的。
4.6 本章小结
嗯，这一章的内容就这些。总结下来就三句话：
- 修饰符：
__global__是入口，__device__是工具，别混用。 - 参数传递：小值传值，大结构传指针，别让寄存器爆炸。
- 异步与同步：默认异步，善用流，只在必要时同步。
这些是核函数编写的地基。地基不稳，后面优化再多也是白搭。下一章我们会深入线程模型，看看 block 和 grid 到底怎么配才能榨干 GPU 的性能。