CUDA学习笔记
文章平均质量分 58
Tonson_
这个作者很懒,什么都没留下…
展开
-
CUDA——设备占用率(occupation)
设备占用率(occupation)1、公式设备占有率是用来衡量核函数的配置分配是否是一个高效率的分配方式。occupation=(active warps)/(max warps)2、max warpsmax warps是每个SM中最多的线程束数量,是一个常量。可以通过查询本机的GPU架构,可知每个SM里最大的可分配的线程数MAX_THREAD_PER_SM(我的机子为2048),以及每个warp包含32个thread。所以,max warps=MAX_THREAD_PER_SM/32。3、a原创 2021-06-19 22:21:21 · 1553 阅读 · 0 评论 -
CUDA——SM中warp调度器调度机制&&访存延迟隐藏
SM中warp调度器调度机制&&访存延迟隐藏核函数中并不是所有线程一起启动执行的,核函数的执行是以线程束(warps)作为单位,warps的执行由warp调度器进行调度,一个调度器只能调度一个warp去执行指令,一个warp里的所有线程几乎是同时执行的。以一个warps调度器为例子:假设一个核函数开启了128个线程,那么其被划分成4个warps。1、调度器调度warp0执行指令。2、warp0挂起(在进行访存),调度器调度warp1执行指令3、以此类推,只要有处于Ready原创 2021-06-19 20:02:33 · 3720 阅读 · 1 评论 -
cublas库中cublasStatus_t枚举变量对应的意思
cublas库中cublasStatus_t枚举变量对应的意思typedef enum{ CUBLAS_STATUS_SUCCESS =0, CUBLAS_STATUS_NOT_INITIALIZED =1, CUBLAS_STATUS_ALLOC_FAILED =3, CUBLAS_STATUS_INVALID_VALUE =7, CUBLAS_STATUS_ARCH_MISMATCH =8, CUBLAS_STATUS_MA原创 2021-06-05 20:48:31 · 1129 阅读 · 0 评论 -
cudaGetErrorString与cudaGetLastError组合运用
cudaGetErrorString与cudaGetLastError组合运用前言在利用集成开发环境(如:visual studio)编写CUDA代码时,编译时是分开编译的,CUDA部分用的是nvcc编译器,这样VS就无法反馈我们CUDA代码的错误,有时候造成了不可避免的内存泄露或者程序出错,自己都不知道,只能一个断点一个断点调试。因此,CUDA也提供了一些函数可以捕捉到error,让我们Debug的时候更方便一点。cudaGetLastError通过翻阅CUDA手册,我们可知宿主线程维护着一个初始原创 2021-06-05 20:20:45 · 4412 阅读 · 0 评论 -
傅里叶变化与卷积和互相关操作的转换
已知有二维信号f(x,y),g(x,y),其对应的傅里叶变化结果为F(x,y),G(x,y)。本篇文章不对公式进行推导,只是运用。有兴趣的话可以自行百度。1、卷积与傅里叶的转换FT{f(x,y)*g(x,y)}=F(x,y)G(x,y)其中,FT{}表示傅里叶操作,也就是两个信号卷积后进行傅里叶变换,结果会是各自傅里叶结果的点乘。那么,假设要求的是两个信号的卷积,则只需对各自进行傅里叶变化再进行点乘,最后再逆傅里叶变化,就可以求得结果。2、互相关与傅里叶的转换FT{f(x,y)⊗g(x,y)原创 2021-04-06 13:04:23 · 2733 阅读 · 0 评论 -
自己编写CUDA常用宏定义
编写CUDA常用宏定义#define一个报错宏我们知道CUDA函数几乎都有一个返回值,它们都是表示CUDA函数执行是否成功的变量。比如runtime API函数的返回值一般是cudaError_tcublas库函数的返回值一般是cublasHandle_t我们就可以根据这些返回值来知道函数执行是否成功,若错误,则通过一些系统宏定义来输出错误的位置在哪一个文件的哪一行下面为宏的内容:#define CHECK_CUDA(errorInfo,cudaSuccInfo){ \ if((err原创 2020-09-16 15:51:40 · 1043 阅读 · 0 评论 -
关于CUDA里的cublas和cufft库的句柄问题
CUDA库的句柄cuda库的句柄是用来切换上下文的,一般要用到库里边的函数,需要传入一个句柄参数。这里以cublas库为例子我们先创建一个句柄cublasHandle_t handle;cublasCreate(&handle);接着,调用cublas库的函数cublasFunction(handle,....)然后我们需要销毁句柄(这步不要省略),特别是用cufft库的时候cublasDestroy(handle);句柄所带来的一些Bug如果句柄没有被销毁的话,那么当重原创 2020-06-22 00:59:39 · 1195 阅读 · 6 评论 -
CUDA项目做科学计算应注意的点
最近在写一个图像匹配并且矫正的算法,主要是用互相关去做,然后求解相关系数最大的位置,再进行图像平移矫正。遇到的科学计算问题在做互相关时,往往在频率上去做互相关,那么两幅图之间就要进行傅里叶变换,然后再点乘。这个时候每个像素位置的实部数值已经达到了e17以上,然后我没注意到这个数值的大小已经超过了float的大小,所以导致后面的数值出错。这个Bug找了我接近一个下午的时间。解决方法将...原创 2020-06-22 00:35:06 · 457 阅读 · 0 评论 -
VS+CUDA项目中用到Eigen库时遇到的问题
项目的编译器和库的版本编译器版本:vs2017cuda版本:10.0问题1:Eigen库的代码无法在.cu文件中通过编译。项目中需要用到最小二乘法的LM算法去求解非线性方程的解,由于之前用Eigen库实现过这个功能,而且特别方便,所以这次也打算使用此库。但是Eigen库的代码无法在.cu里面进行编译,我看了一下报错的文件,发现Eigen库里面有一些源码的文件命名和cuda的源码文件命名相同了,可能是因为这个原因导致的编译问题,nvvc编译器无法识别出。解决方法将利用到Eigen库的代码放到.cp原创 2020-06-22 00:00:02 · 1417 阅读 · 11 评论 -
关于CUDA性能优化的补充
之前写过一篇CUDA性能优化的博客https://blog.csdn.net/weixin_44444450/article/details/104535306上面那篇博客是看完CUDA手册里面关于性能优化的内容后总结的,但是有些地方还是不能很完美的理解。所以现在写一篇博客来补充一下。1.关于全局内存的合并访问要想合并访问全局内存,就必须对地址进行一个对齐。设备能通过单个指令将32位,64位...原创 2020-05-04 16:45:16 · 380 阅读 · 0 评论 -
CUDA归约求和Debug版本结果与Release版本不同的解决方法
CUDA归约求和Debug版本结果与Release版本不同的解决方法近日,在写一个CUDA项目时,里面有一个核函数是要求均值,所以采用归约求和算法,然后再去除以总数求得平均值,在Debug调试版本下,数据正常,但当在Release版本下时,数据却出现了错误。找了很久的错误,最后发现是在归约求和这里出现了错误。原因就是在最后一个warp内的归约求和中,按照原理应该是不用加块内线程同步的,也就是_...原创 2020-04-17 11:44:41 · 737 阅读 · 1 评论 -
做CUDA项目的一些心路历程
做CUDA项目时的一些坑1)要注意数据类型,在进行cudaMemcpy()的时候,要注意数据类型的统一,即使你float的数组里存的是数据是没有小数的整数类型,但当你想要把数据从设备端拷贝回主机端的时候,一定要new一个float类型的数组去存放。2)要注意核函数里,数组的索引一定不要超过范围,否则会出现全是0的情况。特别要注意当一个数组用于存储坐标,另一个数组存放数据。用该存储坐标的数组去索...原创 2020-03-09 12:07:33 · 650 阅读 · 0 评论 -
CUDA——性能优化(总结)
CUDA性能优化策略(总结)1)每个块上的线程数给定每网格的线程总数,设计每块的线程数或网格的块数时应该最大化可用计算资源的利用率。a.块的数目块的数目至少要大于等于你的多处理器(SM)的个数,这样才能充分调动所有多处理器。每个多处理器上的块应存在两个或者以上的活动块。保证在线程同步(块同步)时,有多的活动块可用调用。此时,每个块上的共享内存至多为每个多处理器上共享内存的一半。b.线程...原创 2020-02-27 23:49:50 · 3222 阅读 · 0 评论 -
CUDA——性能优化之循环展开
循环展开(#pragma unroll)循环展开顾名思义就是将循环体展开。全部展开或者展开一部分都可以有效提高性能。以下是一个循环体float sum=0;for(int i=0;i<n;++i){ sum+=a[i];}循环部分展开for(int i=0;i<n;i+=2){ sum+=a[i]+a[i+1];}...原创 2020-02-24 01:21:05 · 6062 阅读 · 0 评论 -
CUDA——性能优化之共享内存
一、共享内存的结构1)什么是共享内存?共享内存是GPU的一种稀缺资源,它位于芯片上,所以共享内存空间要比本地和全局内存空间快得多。对于warp里的所有线程,只要线程之间没有任何存储体冲突(bank conflict),访问共享内存就与访问寄存器一样快。2)什么是存储体(bank)?共享内存被划分为同样大小的、可以同时访问的内存块,名为存储体。在计算能力为1.x的设备上,存储体数为16,在2...原创 2020-02-22 18:25:40 · 1658 阅读 · 0 评论 -
CUDA——性能优化之全局内存
CUDA全局内存的合并访问(个人理解)关于warp指令基础知识1)什么是warp?一个线程warp包括32条线程(我的电脑是1个warp包括32条线程)。它位于多处理器中。2)warp指令发射warp的一个指令,即该warp的32条线程一起执行的该条指令。多处理器会花费 该条指令 个时钟周期。3)控制流指令任何流控制指令( if , switch , do , for , while...原创 2020-02-20 16:49:43 · 1594 阅读 · 0 评论 -
CUDA——流管理
CUDA 流(streams)异步并发执行当使用异步函数时,即实现了设备和宿主之间的并发执行。异步函数:在设备完成请求的任务之前,控制就会返回到应用程序(宿主)。这些函数包括:1.核函数2.执行内存复制并以Async为后缀的函数3.执行 设备<一>设备 内存复制的函数4.设置内存的函数。流的优点应用程序通过流(streams)管理并发。流是一个顺序执行的操作序列。另...原创 2020-02-16 17:02:56 · 426 阅读 · 1 评论 -
CUDA——纹理内存(纹理对象)
纹理内存(二)1.2.2纹理对象API使用纹理对象包括1.纹理对象的创建2.访问纹理内存3.纹理对象销毁纹理对象的创建1)创建纹理对象,cudaCreateTextureObject();cudaCreateTextureObject(cudaTextureObject_t *pTexObject, const struct cudaResourceDesc *pResDesc...原创 2020-02-10 17:32:27 · 1328 阅读 · 0 评论 -
CUDA——纹理内存(纹理参考)
纹理内存1.1纹理类型CUDA支持GPU上的一部分纹理硬件(它们原本是为图形处理而设计的),所以,纹理内存一般会用于加速图片的滤波和图片的resize操作上。纹理内存实质上市全局内存的特殊形态,全局内存被绑定为纹理内存,对其的读,写操作将通过专门的texture cache(纹理缓存)进行。纹理缓存的优势:纹理缓存具备硬件插值特性,可以实现最近邻插值和线性插值。纹理缓存访问二维矩阵的领域会...原创 2020-02-09 15:50:55 · 1145 阅读 · 0 评论