文章目录
参考blibli
P2 CUDA编程入门01-GPU硬件架构综述
bank的访问冲突
冲突的情况
规约的一个有效的算法
p3 CUDA编程模型
CUDA程序执行流程
##### Block和grid的关系
CUDA程序层次结构
CUDA kernel函数的grid, block调用情况
注意用__global__定义的kernel是异步的:调用这个函数cpu就走了,不管GPU是否算完,不会等待GPU完成,而GPU会一个个执行kernel函数
CUDA内置变量
P4向量加法
# GPU申请内存的函数:
cudaMalloc((void **)&dx, nbytes) //你首先得给我一个pointer to pointer of sth, 然后你要申请多少内存空间
P5Grid-Block-Warp-Thread
CUDA程序层次结构
P6GPU内存介绍
左边是Device的图
右边是单个Multi-processor的图
GPU内存类型
P7内存如何管理
内存管理,全局内存是我的数据主要存放的地方,共享内存就是block内线程数据共享和同步
CPU内存
栈的内存有限几M,堆的话看电脑32G,64G
GPU内存
GPU全局内存分配释放
统一(unified)内存分配释放
避免GPU和CPU数据传输,因为都可以访问
同步拷贝: src全部拷贝给dst,这个函数才会终止执行下一个函数
异步拷贝: 只要启动这个函数就不管了(留着GPU与CPU自己完成),就执行这个函数后面的语句,异步拷贝有一个缺点:就是你不知道是否拷贝完成,你就使用他有可能使用的是以前的旧数据
共享内存
常用的使用共享内存的两种方式
P8内存管理,代码示例
申请内存(cudaMalloc/GPU/, malloc或者cudaMallocHost/CPU/)->内存拷贝(cudaMemcpy)->调用核函数 <<<blockIdx,threadIdx>>> ->内存拷贝->释放内存
P9CUDA程序执行与硬件映射
grid->Block->Thread 对应 Device->SM->CUDA Core
P10什么是规约算法
CUDA不支持global sync,因为全局同步可能造成:
- 有block闲置
- deadlock
solution: Decompose into multiple kernels 设计规约算法
多个块block规约如何进行同步:
CUDA没有全局同步的原因:
死锁状态
解决办法:对规约算法分成多个kernel(本来是多个block进行,现在分/decopose多个kernel),相当于在原来的基础上时间纵轴上加了一个里程碑,一个里程碑完成才能进行下一个里程碑,这样block的资源就释放了
调用kernel示意图,即一个规约运算分为两个kernel但是在调用第二个kernel参与运算时,完成第一个kernel运算会到达一个同步点
P11并行规约算法-二叉树算法
策略:
对于相加的规约算法来说,他的示意图如下,每一次线程操作减半:
上述示意图的伪代码:
P12并行规约算法-改进warp divergence
p11的算法,同一warp中的线程执行命令不一样,
P13并行规约算法-改进共享内存访问,消除冲突
P14并行规约算法-改进全局内存访问
P17并行规约算法-成功优化的关键
voliate关键字