CUDA 编程
zkxhlbt
这个作者很懒,什么都没留下…
展开
-
CUDA编程 基础与实践 学习笔记(十)
线程束(warp)一个GPU由多个SM组成,一个SM上可以放多个线程块,不同线程块之间并行或顺序执行。一个线程块分为多个线程束,一个线程束由32个线程(有连续的线程号)组成。从更细粒度来看,一个SM以一个线程束为单位产生、管理、调度、执行线程。福特架构之前,每个warp只有一个程序计数器,需要注意分支发散问题。一些严重的分支发散会极大降低性能。从福特架构开始,引入了独立线程调度机制。书里关于这个感觉没有讲的特别明白,没有太搞清楚和分支发散的逻辑关系。https://baijiahao.baidu.原创 2022-04-24 20:19:36 · 239 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(十四)
CUDA标准库看到了熟悉的cuDNN~原创 2022-04-23 16:38:29 · 854 阅读 · 1 评论 -
CUDA编程 基础与实践 学习笔记(十二)
统一内存管理(某些功能在windows系统下受到限制)好处:使编程更简单,不用手动copy内存。可能提供更好的性能允许超量分配(和虚拟存储类似?)统一内存被当作全局内存来使用的,必须在主机端定义或分配。动态统一内存:静态统一内存:...原创 2022-04-23 16:20:57 · 535 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(九)
原子操作atomicAdd(&addr, val)这个在SNN部分反向传播的cuda代码里面见过,原子操作也在数据库课程里面学过,是不受其他事务影响的“读-改-写”操作。线程的执行不能保证特定顺序,因此需要原子操作来保证变量读写不受影响。所有的原子操作。第一个地址既可以是全局内存,也可以是共享内存。都是device函数,只能在核函数中使用。...原创 2022-04-23 11:44:31 · 349 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(八)
共享内存共享内存可以被程序员直接操控,减少对全局内存的访问,实现高效的线程块内部通信,也可以提高全局内存访问的合并度。数字规约算法。可以并行快速算出数组的和。共享内存限定符__shared__ 变量名可以以s_开头,提醒自己这是共享内存变量。所有线程都有共享内存变量的副本,每个副本不一样但共用一个变量名。核函数对于共享内存的操作作用在所有线程上。因为还没有在实际中用过,所以有一个忧虑的点:其实程序一开始就要将数据从全局变量中移到共享内存中,这一步的耗时如何?如果算术强度不大或者不对共享内存进行多次原创 2022-04-22 00:18:43 · 377 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(七)
条过半…7.1 全局内存的合并和非合并访问从fermi开始,有了SM级别的L1 cache和设备层次的L2 cache. 有了访问全局内存的请求后,先L1,再L2,都不成功的话最后全局内存DRAM中取。一次数据传输在默认情况下是32B.合并度:线程束请求的字节数和该请求导致的数据传输处理的字节数之比。可以视为一种资源利用率的表征。100%则为合并访问,否则为非合并访问。(为啥数据传输处理的字节数会多余请求的字节数,因为要像cpu的访存机制一样从全局取到cache中吗?答:这是基操,不在计算范围内。和访原创 2022-04-20 23:00:39 · 559 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(六)
6.1 CUDA的内存组织内存存在一种层级结构,大容量的一般latency高,小容量的一般latency低。这一点在计算机组成原理中学到过。这里有个疑问,表中的局部内存是放在芯片外部的,这里怎么画在了内部?还是我没有理解清楚?答:局部变量属于全局变量的一部分所以在芯片外。当寄存器放不下了,就会放在局部变量中,由寄存器做判断。6.2 集中不同类型的内存:(1)全局内存,核函数所有线程都能访问,就是显卡的显存,低速高延迟(速度慢是相对于数据处理而言的)。负责host->device, devic原创 2022-04-19 23:17:32 · 971 阅读 · 1 评论 -
CUDA编程 基础与实践 学习笔记(五)
5.1 CUDA事件计时https://blog.csdn.net/qq_24990189/article/details/89602618注意,从这里之后的代码都要用CHECK来检查runtime API的错误。一个例外:cudaEventQuery(),这个函数可能返回的是cudaErrorNotReady,并不代表真的错了。计时的步骤:先定义两个cudaEvent_t变量,并初始化事件,然后开始record。第五行要query一下事件队列。很容易记忆,都是成对的操作。这里既可以是hos原创 2022-04-18 22:22:55 · 613 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(四)
CUDA程序的错误检测这个还蛮关键的,当时看deformable convolution的cuda代码就看到很多CHECK(…),一头雾水。现在学到了感觉还蛮开心的,老朋友了~基本就和自己实现的一个try catch语句块一样。书中给出的模板为:https://github.com/brucefan1983/CUDA-Programming/blob/master/src/04-error-check/error.cuh注意这里的后缀名,是.cuh。错误检测分为runtime API 错误检测和原创 2022-04-17 23:20:07 · 546 阅读 · 0 评论 -
CUDA编程 基础与实践 学习笔记(三)
3.2 CUDA程序的基本框架:cuda runtime api: docs.nvidia.com/cuda/cuda-runtime-api。关于cudaMalloc函数:https://www.cnblogs.com/thkkk/p/14949014.html从引用的角度去理解这个函数更好理解。可以简化地写为double *d_x;cudaError_t cudaMalloc(&d_x, M);cudaError_t cudaFree(void* address);host函数原创 2022-04-17 22:27:15 · 395 阅读 · 1 评论 -
CUDA编程 基础与实践 学习笔记(二)
环境安装过程略2.1-2.2NVCC是CUDA的编译器驱动,将C++代码给C++编译器,例如g++,自己编译剩下的部分。GPU只是一个device,需要gpu来调度。CUDA源文件拓展名为.cu 一个典型的CUDA代码如下:CUDA核函数两个特点:被限定词__global__来修饰必须返回void这两个位置可以混,但一般还是__global__开头。在上述程序中,<<<grid, block>>>是标准写法,grid规定了block的数量,而b原创 2022-04-17 16:35:46 · 940 阅读 · 1 评论 -
CUDA编程 基础与实践 学习笔记(一)
开个坑…用pytorch太不灵活了,学习一下CUDA编程。研究生都第二年了,坚持不下去就有点丢人了奥。书名见标题,樊哲勇老师的著作,清华大学出版社出版。重点看前面十二章。初次阅读,如有理解错误恳请大家批评指正~mua第一章GPU(graphics processing unit),显卡。与CPU的区别:CPU有更多晶体管,用于数据缓存和流程控制,只有少数几个逻辑计算单元,适合完成复杂的逻辑计算;GPU有数千个核心,适合大规模矩阵运算。(GPU的DRAM和CPU的DRAM是通过PCIe总线来原创 2022-04-17 13:54:48 · 2902 阅读 · 2 评论