本篇博客为高性能CUDA应用设计与开发第七章阅读笔记
CUDA上下文环境
cuda上下文环境
囊括了所必需的驱动状态信息,如虚拟地址空间,流,事件,已分配的内存块等
GPU设备与上下文
GPU设备驱动通过设备驱动程序为应用程序提供多个上下文环境
就可以使单个CUDA应用程序使用多个设备
同一时刻只能有一个上下文环境处于活动状态
所以需要操作多个设备时,需要用cudasetDevice()切换上下文环境
cuda在第一次调用一个改变驱动状态的函数时会自动默认创建一个上下文环境
如cudaMalloc()
默认在 GPU 0 上创建上下文
上下文与流
cuda程序通过将操作排队到流中,管理任务和并行
创建上下文时会隐式的创建一个流,从而命令可以在设备中排队等待执行
cudaMemcpy ,kernel函数等都会自动放入默认流中串行执行
除非指定不同的流,kernel函数流的指定如下
kernel<<<nBlocks, nThreadsPerBlocks, 0,stream>>>(para)
第三个可填参数是共享内存大小的分配
当kernel函数之间需要并发执行时要使用多个流
任务并发与同步
显式同步
cudaEventRecord()
我的理解是他是往当前流中异步发送了一个任务(记录时间)
由于同一个流中的顺序执行,使得他在之前的任务完成之后才能将传入的事件标记为完成
cudaStreamSynchronize() //同步流与cpu线程
cudaDeviceSynchronize() //在cpu线程中同步所有流
cudaStreamWaitEvent() //等待某个事件结束后再执行该流上的任务(这个任务在cudaStreamWaitEvent调用之后指定)
cudaStreamQuery() //查询一个流任务是否完成
隐式同步
部分主机操作在其完成前会强制所有流暂停执行
在使用下面这些主机操作时需注意,它们可能会停止所有的并发操作
- 页锁定主机内粗的分配
- 设备内存的分配
- 设备内存设置
- 设备到设备的内存拷贝
- L1缓存与共享内存之间的配置转换
p.s. 我印象中内存的拷贝有异步api的
同一GPU并发执行kernel注意事项
运行多个kernel会改变索引的局部性
增加L2缓存的未命中率,甚至消除缓存的有效性
还会引入其他问题
如bank冲突和内存分区冲突
内存映射
这部分书上主要讲了页锁定内存
感觉没另一本书讲的详细,之前博客写了不复述了