- 博客(32)
- 收藏
- 关注
RSS订阅原创 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
1609
原创 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
4076
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
1172
原创 cudaGetErrorString与cudaGetLastError组合运用
cudaGetErrorString与cudaGetLastError组合运用前言在利用集成开发环境(如：visual studio)编写CUDA代码时，编译时是分开编译的，CUDA部分用的是nvcc编译器，这样VS就无法反馈我们CUDA代码的错误，有时候造成了不可避免的内存泄露或者程序出错，自己都不知道，只能一个断点一个断点调试。因此，CUDA也提供了一些函数可以捕捉到error，让我们Debug的时候更方便一点。cudaGetLastError通过翻阅CUDA手册，我们可知宿主线程维护着一个初始
2021-06-05 20:20:45
4706
原创 傅里叶变化与卷积和互相关操作的转换
已知有二维信号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
2821
原创 C++智能指针：shared_ptr用法详解
C++智能指针：shared_ptr用法详解shared_ptr是C++11里的新特性，其包装了new操作符在堆上分配的动态对象。如：shared_ptr<int> sp1(new int(100));//相当于//int *sp1=new int(100);//auto sp1=make_shared<int>(100);...
2020-12-14 21:13:01
9870
原创 C++ new申请二维动态数组详解
C++ new申请二维动态数组详解new关键字new关键字用于申请一段动态内存空间，位于堆上，在程序运行过程中才执行。可以用来创新数组和实例化对象。如图所示：用new申请一维动态数组，大小为5。int *arrayPtr=new int[5];其中new后面接的是对象，上面为一个int[5]的数组，其返回值为指向数组元素的指针变量。即申请了sizeof(int)*5个字节大小的内存，并返回其地址值存放于指针arrayPtr中。二维数组动态申请方法一int **arrayPtr=new in
2020-12-01 21:02:59
10021
3
原创 避免使用指针变量作为函数参数去申请动态内存
如果函数的参数是指针，避免用该指针去申请动态内存在C程序中，如果函数的参数是一个指针，那么应该坚决避免使用该指针申请动态内存。来看下面一个示例代码：void GetMemory(char *p,int num){ p=(char*)malloc(sizeof(char)*num);}int main(void){ char *str=NULL; GetMemory(str,10); if(str==NULL) { printf("str==NULL\n"); } else {
2020-11-18 11:19:24
238
原创 有符号执行位操作导致的BUG
尽量避免对未知的有符号位执行位操作在C语言中，如果在未知的有符号上执行位操作，很可能会导致缓冲区溢出，从而在某些情况下导致攻击者执行任意代码，同时，还可能会出现出乎意料的行为或编译器定义的行为。如下列代码：#include<stdio.h>int main(void){ int y=0x80000000;//2^31,已超出int类型的范围。 printf("%U",y>>24);//以十进制无符号形式输出。 return 0;}首先定义了一个整型变量y为0x8
2020-11-13 15:54:16
381
原创 谨慎extern申明外部变量
extern关键字用法在C语言程序中，我们用extern关键字对某个变量作“外部变量申明”，表示该变量是一个已经定义的外部变量，编译器就会自动地在所有源文件里面查找该变量地定义。一、在源文件中扩展全局变量的有效范围全局变量的作用范围只限于定义处到文件结束。若想在定义之前使用全局变量，就需要extern关键字来解决。#include<stdio.h>int main(){ int result; extern int number1; extern int number2; r
2020-11-11 23:40:21
1301
原创 C语言typedef关键字用法总结
C语言 typedef用法总结C语言允许我们使用typedef关键字来重新定义类型名，当跨平台移植程序时，我们只需要修改以下typedef的定义即可，而不用对其他源代码做任何修改。用法一：为基本数据类型定义新的类型名typedef int myIntmyInt value //相当于int value重新定义int类型的别名为Tdata用法二：为自定义数据类型(结构体、联合体和枚举类型)定义简洁的类型名称以结构体为例子：struct Complex{ double image; d
2020-11-11 10:07:11
902
原创 C语言--数据溢出导致的缓冲区(堆栈)损坏
防止数据类型变量产生回绕与溢出产生以上大部分问题的原因主要是：1.不同数据类型的算术操作，产生了回绕与溢出。2.不同数据类型的逻辑运算操作，产生了回绕与溢出。3.存储的数据大小超过了数据类型的最大值或最小值，产生了溢出和回绕。一、显示地指出某个变量是signed char类型还是unsigned char来执行算数运算由于不同编译器对char类型的变量定义不同，有可能是signed char，也有可能是unsigned char，所以，只有我们显示地指定数据类型，才能得到唯一的结果。二、使用s
2020-11-10 09:41:54
1912
原创 自己编写CUDA常用宏定义
编写CUDA常用宏定义#define一个报错宏我们知道CUDA函数几乎都有一个返回值，它们都是表示CUDA函数执行是否成功的变量。比如runtime API函数的返回值一般是cudaError_tcublas库函数的返回值一般是cublasHandle_t我们就可以根据这些返回值来知道函数执行是否成功，若错误，则通过一些系统宏定义来输出错误的位置在哪一个文件的哪一行下面为宏的内容：#define CHECK_CUDA(errorInfo,cudaSuccInfo){ \ if((err
2020-09-16 15:51:40
1094
原创 关于CUDA里的cublas和cufft库的句柄问题
CUDA库的句柄cuda库的句柄是用来切换上下文的，一般要用到库里边的函数，需要传入一个句柄参数。这里以cublas库为例子我们先创建一个句柄cublasHandle_t handle;cublasCreate(&handle);接着，调用cublas库的函数cublasFunction(handle,....)然后我们需要销毁句柄(这步不要省略)，特别是用cufft库的时候cublasDestroy(handle);句柄所带来的一些Bug如果句柄没有被销毁的话，那么当重
2020-06-22 00:59:39
1238
6
原创 CUDA项目做科学计算应注意的点
最近在写一个图像匹配并且矫正的算法，主要是用互相关去做，然后求解相关系数最大的位置，再进行图像平移矫正。遇到的科学计算问题在做互相关时，往往在频率上去做互相关，那么两幅图之间就要进行傅里叶变换，然后再点乘。这个时候每个像素位置的实部数值已经达到了e17以上，然后我没注意到这个数值的大小已经超过了float的大小，所以导致后面的数值出错。这个Bug找了我接近一个下午的时间。解决方法将...
2020-06-22 00:35:06
475
原创 VS+CUDA项目中用到Eigen库时遇到的问题
项目的编译器和库的版本编译器版本：vs2017cuda版本：10.0问题1：Eigen库的代码无法在.cu文件中通过编译。项目中需要用到最小二乘法的LM算法去求解非线性方程的解，由于之前用Eigen库实现过这个功能，而且特别方便，所以这次也打算使用此库。但是Eigen库的代码无法在.cu里面进行编译，我看了一下报错的文件，发现Eigen库里面有一些源码的文件命名和cuda的源码文件命名相同了，可能是因为这个原因导致的编译问题，nvvc编译器无法识别出。解决方法将利用到Eigen库的代码放到.cp
2020-06-22 00:00:02
1482
11
原创 C++排序算法(二路归并排序)
C++排序算法(二路归并排序)归并排序算法是一个可以在最坏情况下依然保持O(nlogn)运行时间的确定性排序算法。归并排序的主体结构属典型的分治策略，我们可以先进行无序向量的递归分解，再进行有序向量的逐层归并。代码实现#include<iostream>using namespace std;template<class T>void merge(T *array, int left,int mi,int right){ T* A=array+left;
2020-05-13 12:02:22
1376
原创 C++ 查找中位数
查找C++无序数列中位数1.将无序数列转化为有序数列，然后直接通过索引找到中位数2.已知中位数 要求左边小于（大于）中位数，右边大于（小于）中位数，且左边数列大小等于右边数列大小 。 而对左边数列或者右边数列数值的排列顺序无要求。所以，我们可以采用优先队列进行解决。代码实现#include<iostream>#include<queue>//优先队列#include<vector>#include<assert.h>#include<l
2020-05-11 18:08:17
5345
原创 C++ 二分查找实现
C++ 二分查找二分查找法是针对于有序向量中，查找该向量中某一个数。二分查找法是通过将向量中点的值与要查找的值进行比较，然后每次减半去缩小查找范围。代码实现#include<iostream>using namespace std;//时间复杂度O(log)template<class T>int binaryFind(T *arr,const T &value, int left,int right){ //若数组中value有重复，则会找到
2020-05-11 17:33:16
291
原创 C++ 广度优先搜索(邻接表+队列实现)
C++ 广度优先搜索(邻接表+队列实现)广度优先搜索的结果是图的生成树，也就是极小连通图（n个顶点，n-1条边）广度优先搜索 相当于 树结构的层次遍历，因此我们可以套用树的层次遍历，利用队列去解决图的广度优先搜索...
2020-05-11 16:18:11
1142
原创 C++ 深度优先搜索(邻接表+栈实现)
C++ 深度优先搜索(邻接表+栈实现)深度优先搜索的结果是图的生成树，也就是极小连通图（n个顶点有n-1条边）图的生成树的结果是不唯一的，它是根据你的存储方式的变化而变化的。深度搜索步骤：1.确定一个开始顶点，并将该顶点标记为已访问。2.访问开始顶点的第一个邻接顶点，if(未被访问)将该邻接顶点标记为已访问。将该邻接顶点作为新的开始顶点。重复2else搜索下一个邻接顶点。重复23.当开始顶点的所有邻接顶点均已被访问回溯到上一个开始顶点，重复24.直到所有的顶点都已被访问过(没有点
2020-05-11 12:30:59
2160
原创 关于CUDA性能优化的补充
之前写过一篇CUDA性能优化的博客https://blog.csdn.net/weixin_44444450/article/details/104535306上面那篇博客是看完CUDA手册里面关于性能优化的内容后总结的，但是有些地方还是不能很完美的理解。所以现在写一篇博客来补充一下。1.关于全局内存的合并访问要想合并访问全局内存，就必须对地址进行一个对齐。设备能通过单个指令将32位，64位...
2020-05-04 16:45:16
400
原创 CUDA归约求和Debug版本结果与Release版本不同的解决方法
CUDA归约求和Debug版本结果与Release版本不同的解决方法近日，在写一个CUDA项目时，里面有一个核函数是要求均值，所以采用归约求和算法，然后再去除以总数求得平均值，在Debug调试版本下，数据正常，但当在Release版本下时，数据却出现了错误。找了很久的错误，最后发现是在归约求和这里出现了错误。原因就是在最后一个warp内的归约求和中，按照原理应该是不用加块内线程同步的，也就是_...
2020-04-17 11:44:41
758
1
原创 梯度及梯度下降法
梯度梯度是哈密尔顿算子（Nabla算子）直接作用于函数F( r )（不论F是矢量还是标量），当F是标量时，梯度即为矢量。这里呢，解释一下什么是哈密尔顿算子实际上，哈密尔顿算子就是一个微分算子，就是我们熟悉的倒三角符号∇\nabla∇。标量场中的梯度指向标量场中增长最快的方向，梯度的长度就是这个最大的变化率。...
2020-03-22 23:00:50
715
原创 做CUDA项目的一些心路历程
做CUDA项目时的一些坑1）要注意数据类型，在进行cudaMemcpy()的时候，要注意数据类型的统一，即使你float的数组里存的是数据是没有小数的整数类型，但当你想要把数据从设备端拷贝回主机端的时候，一定要new一个float类型的数组去存放。2）要注意核函数里，数组的索引一定不要超过范围，否则会出现全是0的情况。特别要注意当一个数组用于存储坐标，另一个数组存放数据。用该存储坐标的数组去索...
2020-03-09 12:07:33
656
原创 CUDA——性能优化（总结）
CUDA性能优化策略（总结）1）每个块上的线程数给定每网格的线程总数，设计每块的线程数或网格的块数时应该最大化可用计算资源的利用率。a.块的数目块的数目至少要大于等于你的多处理器（SM）的个数，这样才能充分调动所有多处理器。每个多处理器上的块应存在两个或者以上的活动块。保证在线程同步（块同步）时，有多的活动块可用调用。此时，每个块上的共享内存至多为每个多处理器上共享内存的一半。b.线程...
2020-02-27 23:49:50
3353
原创 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
6368
原创 CUDA——性能优化之共享内存
一、共享内存的结构1)什么是共享内存？共享内存是GPU的一种稀缺资源，它位于芯片上，所以共享内存空间要比本地和全局内存空间快得多。对于warp里的所有线程，只要线程之间没有任何存储体冲突(bank conflict)，访问共享内存就与访问寄存器一样快。2)什么是存储体（bank）？共享内存被划分为同样大小的、可以同时访问的内存块，名为存储体。在计算能力为1.x的设备上，存储体数为16，在2...
2020-02-22 18:25:40
1790
原创 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
1643
原创 CUDA——流管理
CUDA 流(streams)异步并发执行当使用异步函数时，即实现了设备和宿主之间的并发执行。异步函数：在设备完成请求的任务之前，控制就会返回到应用程序（宿主）。这些函数包括：1.核函数2.执行内存复制并以Async为后缀的函数3.执行 设备<一>设备 内存复制的函数4.设置内存的函数。流的优点应用程序通过流(streams)管理并发。流是一个顺序执行的操作序列。另...
2020-02-16 17:02:56
444
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
1391
原创 CUDA——纹理内存（纹理参考）
纹理内存1.1纹理类型CUDA支持GPU上的一部分纹理硬件（它们原本是为图形处理而设计的），所以，纹理内存一般会用于加速图片的滤波和图片的resize操作上。纹理内存实质上市全局内存的特殊形态，全局内存被绑定为纹理内存，对其的读，写操作将通过专门的texture cache（纹理缓存）进行。纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性插值。纹理缓存访问二维矩阵的领域会...
2020-02-09 15:50:55
1210
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人