写在前面:
- NVIDIA CUDA初级教程笔记 p11~p12
- 连接:https://www.bilibili.com/video/BV1kx411m7Fk?p=12
有效的数据并行算法 + 针对GPU架构特性的优化 = 最优性能
文章目录
1.Parallel Reduction 并行规约
这个过程类似篮球锦标赛的淘汰过程:n个元素进行log(n)个回合,如何在CUDA上实现?
//累加存在shared memory内的元素
__shared__ float partialSum[element_num];
unsigned int t = threadIdx.x;
//步长:1,2,4...
for (unsigned int stride = 1; stride < blockDim.x; stride*=2)
{
__syncthreads();//保证每一步做完之后再进行下一步
//在同一块shared memory里面进行累加
//当步长增加时,多余的线程在干什么?没事干
if (t%(2*stride)==0)
{
partialSum[t] += partialSum[t + stride];
}
}
如果我们改进这个过程会怎么样?
每一轮所需要的线程数依然是减半的,但是线程所处的位置不同
第二种可以将提前完成的线程的硬件资源释放,用来做其他的事情。
元素求和源代码
#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdlib.h>
/*
*处理有N个元素的并行规约
*/
# define N 2048
# define threadsPerBlock 512
# define blocksPerGrid (N+threadsPerBlock-1)/threadsPerBlock//4
__global__ void ReductioinSum_0(float* d_a, float* d_partial_sum)
{
//申请共享内存, 存在于每个block中
__shared__ float partialSum[threadsPerBlock];
//确定索引
unsigned int i = threadIdx.x+blockIdx.x*blockDim.x;
unsigned int tid = threadIdx.x;
//传global mem数据到shared memory
partialSum[tid] = d_a[i];
//在共享存储器中进行规约
//步长:1,2,4...
for (unsigned int stride = 1; stride < blockDim.x; stride*=2)
{
__syncthreads();//保证每一步做完之后再进行下一步
//在同一块shared memory里面进行累加
//当步长增加时,多余的线程在干什么?没事干
if (tid%(2*stride)==0)
{
partialSum[tid] += partialSum[tid + stride];
}
}
//将当前block的计算结果写回输出数组
if (tid == 0)
d_partial_sum[blockIdx.x] = partialSum[0];
}
__global__ void ReductioinSum_01(float* d_a, float* d_partial_sum)
{
//申请共享内存, 存在于每个block中
__shared__ float partialSum[threadsPerBlock];
//确定索引
unsigned int i = threadIdx.x + blockIdx.x*blockDim.x;
unsigned int tid = threadIdx.x;
//传global mem数据到shared memory
partialSum[tid] = d_a[i];
//在共享存储器中进行规约
//步长:4.2.1.
for (unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2)
{
__syncthreads();
if (tid<stride)
{
partialSum[tid] += partialSum[tid + stride];
}
}
//将当前block的计算结果写回输出数组
if (tid == 0)
d_partial_sum[blockIdx.x] = partialSum[0];
}
int main(void)
{
float a[N],partial_sum[blocksPerGrid];
//初始化示例数据
for (int i = 0; i<N; i++)
{
a[i] = i;
}
int size = N*sizeof(float);
//分配显存空间
float *d_a,*d_partial_sum;
cudaMalloc((void**)&d_a, size);
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMalloc((void**)&d_partial_sum, blocksPerGrid*sizeof(float));
ReductioinSum_01 <<<blocksPerGrid, threadsPerBlock >> >(d_a, d_partial_sum);
cudaMemcpy(partial_sum, d_partial_sum,blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost);
float sum = 0;
for (int i = 0; i < blocksPerGrid; i++) {
sum += partial_sum[i];
}
printf("all sum = %.2f", sum);
cudaFree(d_a);
cudaFree(d_partial_sum);
return 0;
}
2. Warp分割
Warp分割:块内线程如何划分wrap
通晓warp分割有助于:减少分支发散,让warp尽早完工
- Block被划分为以32为单位的线程组,叫做warp
- warp是最基本的调度单元
- Warp里的线程一直执行相同的指令(SIMT)
- 每个线程只能执行自己的代码路径
- Fermi SM有2个warp调度器(Tesla has 1)
- warp里设备切换没有时间代价
- 许多warps在一起可以隐藏访存延时
Warp分割的原则是:threadIdx连续增加的一组
以行为主元的情况:
warp分支分散:会降低性能
例如:给定warpSize=32,以下代码是否有哪个warp存在分支发散
if(threadIdx.x>15)
{
// 存在 ×
}
if(threadIdx.x>warpSize-1)
{
// 不存在 √
}
第二种更好
在第一轮,右边的warp2和warp3可以腾出来做其他的事情
在第2轮,右边的wap1 warp2和warp3可以腾出来做其他的事情
3.Memory Coalesing 访存合并
上文了解了整个线程的调度和warp的切分后,我们来关注存储优化问题。
CPU-GPU 数据传输最小化
- Host <- - > device数据传输带宽远低于global memory
- 减少传输
- 中间数据直接在GPU分配,操作,释放
- 有时更适合在GPU进行重复运算
- 如果没有减少数据传输的话,将CPU代码一直到GPU可能无法提升性能
- 组团传输
- 大块传输好于小块:10微妙延迟,8GB/s => 如果数据小于80KB,性能将受延迟支配
- 内存传输与计算时间重叠
- 双缓存
Coalesing合并
Global memory 延时:400 ~ 800 cycles
最重要的影响因子!
在Fermi,global memory默认缓存于 一级缓存L1
通过给nvcc指令设置参数 ”-Xptxas -dlcm=cg“可以绕过一级缓存L1:只缓存于二级缓存L2
如果缓存:warp的读写请求落到L1 cache line,只需一次传输
#transaction = #L1 line accessed
如果没有缓存:有一些合并原则
但是传输大小可以减至32字节块
Memory Coalescing 合并访存
“相邻的人搬相邻的砖哈哈哈哈”
合并举例
- 小型kernel拷贝数据时的有效带宽
- 偏移和步长对性能的影响
- 2款GPUs
- GTX 280:compute capability 1.3;峰值带宽141GB/2
- FX 5600:compute capability 1.0;峰值带宽77GB/s
偏移量的影响
步长的影响
Shared memory
- 比global memory快上百倍
- 可以通过缓存数据减少global memory访存次数
- 线程可以同通过shared memory协作
- 用来避免不满足合并条件的访存
- 读入shared memory重排顺序,从而支持合并寻址
4.Bank 冲突
按行读,按列写,读写总有一个是不合并的
读和写都实现合并访存
Texture 纹理
总结
如果遵行一些简单的原则,GPU硬件在数据可并行计算问题上,可以达到很好的性能:
- 有效利用并行性
- 尽可能合并内存访问
- 利用shared memory
- 开发其他存储空间
- Texture
- Constant
- 减少bank冲突
5.SM资源动态分割
Performance Cliff: 增加资源用量后导致并行性急剧下降。例如,增加寄存器数量,除非为了隐藏global memory访存延迟
kernel 启动参数配置介绍
Grid Size试探法
Occupancy 占用率
6.数据预读
用来隐藏访存延时的一个小手段
float m = Md[i];//read global memory
float f = a*b+c*d;//执行指令,不依赖于读内存的操作
float f2 = m*f;//在上一行执行足够多的warp隐藏访存延时以后,再使用global memory
在上图中
load next tile into register //为循环的下一次迭代预读
accumulate dot product //这些指令被足够多的线程执行,从而隐藏了预读内存产生的延时
7.指令混合
指令吞吐量优化
8.循环展开
缺点:可扩展性差,blocksize变化会带来影响。