1. 不能递归
2. 不能有返回值
3. 参数的数量（或者说参数的总字节数）有上限
4. 内部不能声明静态变量

__syncthreads();

同步范围是整个Block，一个线程运行到syncthreads后，就会等待所有其他线程也运行到syncthreads
用另一个方式解释就是，A线程运行到了第N个syncthreads，那么它就会等待其他线程也执行到第N个syncthreads，然后再恢复执行，直到第N+1个syncthreads

myKernel<<<gridDim, blockDim, sharedMemSize, cudaStream>>>(input);

如何确定最佳的 gridDim 和 blockDim 呢？
需要对CUDA Core的调度方式有一定的理解，参照硬件细节
使用cudaOccupancyMaxActiveBlocksPerMultiprocessor / cudaOccupancyMaxPotentialBlockSize等函数也可以自动估算？

// Generated by AI
// 使用 occupancy API 自动帮我们计算一个合理的 block 大小
int minGridSize = 0;  // 返回值，建议的最小 grid 大小（block 数量）
int blockSize = 0;    // 返回值，建议的每个 block 的线程数
 
cudaOccupancyMaxPotentialBlockSize(
    &minGridSize,    // 输出：最小推荐 grid 大小（block 数）
    &blockSize,      // 输出：推荐 block 大小（每个 block 的线程数）
    myKernel,        // 目标 kernel 函数指针
    0,               // 动态分配 shared memory 的字节数（这里是 0）
    0                // 设备编号（通常填 0）
);
 
// 根据推荐的 blockSize 计算我们需要的 blocks 数量，保证覆盖所有数据
int gridSize = (N + blockSize - 1) / blockSize;
 
std::cout << "Recommended block size: " << blockSize << std::endl;
std::cout << "Recommended grid size (number of blocks): " << gridSize << std::endl;
 
// 启动 kernel
myKernel<<<gridSize, blockSize>>>(d_data, N);

// gridDim, blockIdx; blockDim, threadIdx
__global__ void add(slice<const double> a1, slice<const double> a2, slice<double> res) {
    int32_t tid      = blockIdx.x * blockDim.x + threadIdx.x;
    int32_t nthreads = gridDim.x * blockDim.x;
    for(int32_t i = tid; i < res.size(); i += nthreads) {
        res(i) = a1(i) + a2(i);
    }
}

多个 stream = 并发执行的多个 kernel，这被称为 kernel concurrency。现代显卡一般都支持，但也有一些老显卡不支持