Cuda编程：核函数-CSDN博客

本文链接：https://blog.csdn.net/WaitingHereFor/article/details/147961295

CUDA编程中的核函数（Kernel Function）是运行在NVIDIA GPU上的并行计算函数。它是CUDA编程模型的核心，允许开发者通过定义一个核函数来指定要并行执行的任务编程。

1.基本概念

核函数是：

在主机(CPU)代码中调用
在设备(GPU)上执行
由大量线程并行执行的函数

2. 核函数定义及调用

核函数使用__global__修饰符声明：

__global__ void myKernel(int *data, int value) {
    // 核函数代码
}

核函数使用特殊的语法调用，指定执行配置（网格和线程块维度）：

myKernel<<<gridDim, blockDim, sharedMemSize, stream>>>(arguments);

myKernel<<<gridDim, blockDim>>>(deviceData, 5);

执行配置指定了如何组织线程：

gridDim：网格维度（block的数量）
blockDim：线程块维度（每个block中thread的数量）
sharedMemSize：动态共享内存大小（字节）
stream：执行流（默认为0）

<<<128, 256>>> // 128个block，每个block有256个thread

CUDA还支持动态并行，即从核函数中启动其他核函数：

__global__ void childKernel() { /* ... */ }

__global__ void parentKernel() {
    if (threadIdx.x == 0) {
        childKernel<<<1, 32>>>();
    }
}

3. 线程索引

在核函数中，可以使用以下内置变量获取线程索引：

blockIdx.x, blockIdx.y, blockIdx.z：block在grid中的索引
threadIdx.x, threadIdx.y, threadIdx.z：thread在block中的索引
blockDim.x, blockDim.y, blockDim.z：block的维度
gridDim.x, gridDim.y, gridDim.z：grid的维度

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

// 调用
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

4. 核函数限制

不能有返回值（必须返回void）

不能是类的成员函数

不能使用可变数量的参数

不能使用静态变量

不能递归调用

核函数的执行时间通常较长，应尽量避免在核函数中进行大量的串行计算

参数传递有一些限制：

参数通过常量内存传递给设备

参数总大小有限制（通常为4KB）

最好传递指针（设备内存指针）而不是大型数据结构

5. 核函数优化技巧

最大化并行性：使用足够的线程
优化内存访问：合并内存访问
避免线程发散：尽量减少条件分支
合理使用共享内存
隐藏延迟：足够的线程块以充分利用GPU

优化CUDA核函数性能需要从多个维度考虑，以下是核心优化策略：

5.1. 最大化并行利用率

网格与线程块配置优化

合理设置block大小：通常选择128-256个线程/block（经验值为128或256的倍数）
避免过小的grid：确保有足够多的block以充分利用GPU所有SM（流式多处理器）
三维配置：对于图像处理等应用，使用2D或3D线程布局更自然

// 示例：2D线程布局
dim3 blockDim(16, 16);  // 256 threads per block
dim3 gridDim((width+15)/16, (height+15)/16);
kernel<<<gridDim, blockDim>>>(...);

5.2. 内存访问优化

全局内存优化

合并访问（Coalesced Access）：确保连续线程访问连续内存地址
对齐访问：内存地址对齐到32字节/128字节边界
利用L2缓存：CUDA 11+支持设置持久化L2缓存

// 不良模式（非合并访问）
__global__ void badAccess(float *data) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    data[tid * 2] = ...;  // 跨步访问，效率低
}

// 优化模式（合并访问）
__global__ void goodAccess(float *data) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    data[tid] = ...;  // 连续访问
}

共享内存优化

Bank冲突避免：确保线程访问不同bank（32-bank架构）
用于数据复用：适合需要多次访问的数据

__global__ void sharedMemExample(float *input, float *output) {
    __shared__ float sData[256];
    int tid = threadIdx.x;
    
    // 从全局内存加载到共享内存
    sData[tid] = input[blockIdx.x * blockDim.x + tid];
    __syncthreads();
    
    // 使用共享内存进行计算
    output[blockIdx.x * blockDim.x + tid] = sData[tid] * 2;
}

5.3. 计算优化

指令级优化

使用快速数学函数：__expf(), __sinf()等内建函数
避免双精度运算：除非必要，使用单精度(float)
循环展开：减少分支预测开销

// 手动循环展开示例
#pragma unroll 4
for(int i=0; i<4; i++) {
    // 计算代码
}

避免线程发散

最小化条件分支：同一warp内的线程应执行相同路径
重构算法：将发散操作移到核函数外部

// 不良模式（线程发散）
__global__ void divergentKernel(int *data) {
    if(threadIdx.x % 2 == 0) {
        // 路径A
    } else {
        // 路径B - 同一warp内的线程会发散
    }
}

5.4. 资源平衡

寄存器使用优化

减少寄存器压力：使用__launch_bounds__限定符
权衡使用局部内存：当寄存器不足时编译器会自动使用局部内存

// 限制寄存器使用
__global__ __launch_bounds__(256, 4) 
void optimizedKernel(...) {
    // 核函数代码
}

隐藏延迟

足够的并行度：通常需要至少3-5个活跃的block/SM
异步操作：与主机通信使用异步函数

5.5. 高级优化技术

warp级编程

使用warp原语：__shfl_sync, __reduce_add_sync等
warp同步操作：减少__syncthreads()使用

// warp缩减求和示例
__global__ void warpReduce(int *data) {
    int tid = threadIdx.x;
    int val = data[tid];
    
    for(int offset=16; offset>0; offset/=2)
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    
    if(tid == 0) data[0] = val;
}