FineKernels

练习实现最优经典算子。

Softmax

实现概述

实现了两种Softmax CUDA kernel：基础实现和优化实现。

基础实现 (softmax_forward_naive)

线程与Block分配

• Block分配：每个Block负责计算一个样本（row）
• Thread分配：每个Block使用128个线程
• Warp分配：128线程 = 4个Warp（每个Warp 32线程）

内存使用

• 全局内存：输入和输出数据
• 共享内存：存储Warp级中间结果，大小为 num_wraps * sizeof(float)

优化实现 (softmax_forward_optimize)

线程与Block分配

• Block分配：每个Block负责计算4个样本（rows_per_block = 4）
• Thread分配：每个Block使用128个线程
• Warp分配：4个Warp，每个Warp负责一个样本（row）

关键优化

1. 向量化读取：使用float4类型一次读取4个元素，提高内存带宽利用率
2. 寄存器优化：
   • 使用float4 reg_val_f4[float4_per_thread]存储中间结果
   • 减少内存访问次数
3. 内存访问模式：
   • 优化全局内存读取模式，提高缓存命中率
   • 避免非对齐访问

性能对比(N = 65536)

特征维度	基础实现	优化实现	加速比
128	0.0463 ms	0.0185 ms	2.50x
1024	0.5891 ms	0.5914 ms	1.00x

Reduce

实现概述

实现了两种Reduce CUDA kernel：基础实现和优化实现。

基础实现 (reduce_naive)

线程与Block分配

• Block分配：每个Block处理 blockDim.x * 2 个元素
• Thread分配：每个Block使用1024个线程
• Warp分配：1024线程 = 32个Warp（每个Warp 32线程）

实现步骤

1. 数据加载：每个线程加载2个元素并求和
2. Warp级规约：使用 __shfl_down_sync 进行Warp内规约
3. 共享内存存储：将每个Warp的结果存储到共享内存
4. Block级规约：由第一个Warp计算所有Warp结果的总和
5. 结果输出：线程0将最终结果写入输出数组

内存使用

• 全局内存：输入和输出数据
• 共享内存：存储Warp级中间结果，大小为 num_wraps * sizeof(float)

优化实现 (reduce_optimize)

1. 向量化读取：使用float4类型一次读取4个元素，提高内存带宽利用率
2. 原子操作：使用atomicAdd直接累加结果到输出，避免二次规约
3. 内存访问模式：优化全局内存读取模式，提高缓存命中率
4. 减少启动次数：从2次kernel启动减少到1次，降低启动开销

性能对比

元素数量	基础实现	优化实现	加速比	基础带宽	优化带宽
1,048,576	0.1126 ms	0.0154 ms	7.31x	37.24 GB/s	273.07 GB/s
262,144	0.0051 ms	0.0041 ms	1.24x	204.80 GB/s	256.00 GB/s

FastDiagramKernel

实现概述

实现了高效的直方图统计CUDA kernel，用于统计输入数据的频数分布。

功能描述

直方图统计算子用于计算输入数据中各个值的出现频数。本实现针对uint8_t类型数据（0-255范围），统计256个bin的频数分布。

实现策略

线程与Block分配

• Thread分配：每个Block使用256个线程
• Warp分配：256线程 = 8个Warp（每个Warp 32线程）

关键优化

1. 共享内存局部直方图：每个Block在共享内存里维护自己的局部直方图，减少全局内存原子操作冲突
2. 向量化读取：使用uint32_t类型一次读取4个字节，提高内存访问效率

Transpose Kernel

实现概述

实现了四种矩阵转置CUDA kernel，从基础实现到逐步优化的版本，展示了共享内存、缓存优化、Bank冲突消除和向量化预取等关键技术。

四种算子实现及优化原理

1. Naive Transpose (基础实现)

优化原理：直接全局内存访问

• 线程映射：每个线程处理一个矩阵元素
• 内存访问模式：
  • 读取时：连续访问（coalesced）
  • 写入时：非连续访问（strided），导致性能瓶颈
• 性能瓶颈：转置写入操作导致全局内存访问不连续，带宽利用率低

2. Smem Direct Map (共享内存直接映射)

优化原理：共享内存缓存 + 直接映射

• 缓存策略：使用共享内存作为中间缓存，将全局内存访问转换为共享内存访问
• 线程映射：
  • 读取阶段：线程(tx, ty)读取全局内存连续数据到共享内存cache[ty][tx]
  • 写入阶段：通过线性索引计算重新映射到cache[tx][ty]
• Bank冲突问题：直接映射会导致写入时的Bank冲突，多个线程同时访问同一Bank

3. Smem Conflict Free (+1) (Bank冲突消除)

优化原理：共享内存Bank冲突消除

• 关键改进：将共享内存数组声明为cache[SM][SN+1]
• Bank冲突消除原理：
  • GPU共享内存采用32个Bank的组织方式
  • 相邻元素在32的倍数位置会映射到同一Bank
  • 通过增加1列填充，使得同一列的元素映射到不同Bank
• 性能提升：消除Bank冲突后，共享内存访问带宽得到充分利用

4. Smem Double Fetch (Vectorized) (向量化预取)

优化原理：向量化预取 + 线程复用

• 线程复用：每个线程处理2个元素，Block维度减半
• 预取策略：
  • 使用float类型一次读取2个连续元素
  • 减少全局内存事务数量
  • 提高内存带宽利用率
• 共享内存优化：使用一维共享内存数组cache[SM*(SN+1)]
• 向量化写入：写入时也采用向量化方式，提高写入效率

性能对比 (矩阵尺寸: 8192×8192)

实现方案	执行时间	相对Naive加速比	优化技术
Naive Transpose	2.0154 ms	1.00x	基础实现
Smem Direct Map	0.6605 ms	3.05x	共享内存缓存
Smem Conflict Free (+1)	0.6574 ms	3.07x	Bank冲突消除
Smem Double Fetch (Vectorized)	0.6553 ms	3.08x	向量化预取