CUDA学习
CUDA学习
武泗海
石油物探 数值模拟 信号处理
展开
-
CUDA 二维矩阵操作
二维矩阵代码如下:#include #include #include"cuda_runtime.h"#include"device_launch_parameters.h"#include"iostream"using namespace std;/**********************************************************原创 2016-05-05 01:20:43 · 1587 阅读 · 0 评论 -
CUDA耗时Kernel函数导致GPU崩溃解决方案
在CUDA的GPU编程中,通常将大数据分配由众多线程解决,这样每个线程的工作量通常很小,秒秒钟解决。当时,当GPU希望承担CPU一样的耗时线程操作时,经常会出现“驱动已恢复”的崩溃现象。这是由于window系统Tdr所导致的,显卡是用来桌面显示的,所以不允许显卡长期(>2s)被占用,需要把GPU抢回来用于图形显示。因此就出现了上述情况。 解决方案:在Nvidia安装的Nsight monito原创 2016-07-02 22:14:16 · 5065 阅读 · 0 评论 -
CUDA获取错误信息
cudaError_t cudastatus;cudastatus = cudaMalloc((void**)&d_data,NByte);if (cudastatus != cudaSuccess) {qDebug("%s", cudaGetErrorString(cudastatus)); }原创 2017-12-13 20:25:55 · 3383 阅读 · 0 评论 -
CUDA并行规约(相邻配对)
#include "cuda_runtime.h"#include "device_launch_parameters.h"#include #include "math.h"#include "stdlib.h"//错误检查的宏定义#define CHECK(call) \{ \ const cudaError_t status=cal原创 2017-12-17 11:16:39 · 610 阅读 · 0 评论 -
CUDA并行规约(相邻配对-优化)
前文CUDA的并行规约算法的示意图如下,分析可知,相邻之间的线程执行不同的路径,存在线程束分化。为了使得线程束不存在分化,每个warp(32个线程)执行同一指令,可调整相邻的线程的数组索引实现优化。示意图如下图所示,数组的存储位置没变,只是没个线程执行的数组发生了变化,这样的处理模式可以降低相邻线程分化降低,尽早释放后面的线程。实验在GTX1050Ti进行,原创 2017-12-17 12:27:35 · 688 阅读 · 1 评论 -
CUDA并行规约(交错配对)
按照前文多线程的交错配对方式实现并行规约求和方式,实现CUDA版本的并行规约求和,由于这种方式的规约可以避免线程束的分化,因此不需要进行类似于相邻配对那种方式的优化。交错与优化相邻模式相比,计算效率提升到1.14倍,性能提高有限,这主要受限于全局内存的加载和存储模式。并行规约的示意图: #include "cuda_runtime.h"#include "device_lau原创 2017-12-17 13:47:38 · 875 阅读 · 0 评论 -
CUDA并行规约(交错配对-展开规约)
之前的规约计算,每个线程块负责一个对应的数据块。现在将每个线程块负责两个数据块的规约,从而达到消除指令消耗,增加更多的独立指令的调度,来提高性能。下面是展开因子为2的示意图,将相邻的数据块加到当前线程块对应的数据块,然后进行规约求和。#include "cuda_runtime.h"#include "device_launch_parameters.h"#include #inclu原创 2017-12-17 19:53:01 · 830 阅读 · 0 评论 -
CUDA并行规约(交错配对-展开线程)
当线程执行规约函数的循环时,当只剩下一个线程束(warp=32,),甚至更少的时候,可以将后续的循环展开,及32 16 8 4 2 1的循环直接展开书写。这样处理可以避免循环控制和线程同步逻辑。#include "cuda_runtime.h"#include "device_launch_parameters.h"#include #include "math.h"#i原创 2017-12-17 20:59:57 · 814 阅读 · 1 评论 -
CUDA并行规约(交错配对-完全展开-终极版)
通过前面的文章,我们对32以下的迭代循环进行了展开处理,实际上,由于线程块的长度限制(GTX1050Ti是1024),可以说循环次数是确定的,因此可以将循环完全展开,即 1024 512 256 128 64 都可以展开计算,唯一需要注意的每次计算之后都要进行同步,原因在前文已有解释。下面给出代码。#include "cuda_runtime.h"#include "device_lau原创 2017-12-17 23:27:59 · 1230 阅读 · 0 评论 -
CUDA 规约计算性能对比
通过前文,已经把整个规约计算的过程给出了。这里是工程文件,开发环境win1064bit+VS2017+CUDA9.0+1050Ti :http://download.csdn.net/download/qq_17239003/10162546给出整体性能的对比,原创 2017-12-17 23:41:11 · 736 阅读 · 0 评论 -
CUDA动态并行
Nvidia的GPU,设备计算能力>3.5的设备支持动态并行。可实现GPU线程启动核函数,而不用CPU端调用核函数。这在递归调用很有用!在项目要强制生成可重定位的设备代码,win操作:1.指定代码生成:设备项目-属性-CUDA C/C++ -Device- Code Generation 修改为 compute_35,sm_35。 具体数值可根据自身设备修改,但至少是SM_35。2原创 2017-12-18 21:56:08 · 1992 阅读 · 2 评论 -
CUDA-全局内存的区别与用法
CUDA编程中全局内存分为分页内存,固定内存,零拷贝内存,统一虚拟寻址,统一内存地址。一.分页内存 利用cudaMalloc()申请的主机内存,即可分页内存。 特点:可分页内存传输数据到设备时,首先需要分配固定内存,再传递到设备端。float* data;cudaMalloc((void**)&data,sizeof(float)*N);原创 2018-01-02 21:01:40 · 2320 阅读 · 0 评论 -
CUDA 同步和异步
同步操作:主机向设备提交任务,主机将阻塞,直到设备将所提交任务完成,并将控制权交回主机。然后继续执行主机的程序。异步操作:主机向设备提交任务,设备直接开始执行任务,但主机将不再阻塞,而是直接继续执行主机的程序。主机并不会等待设备执行任务完毕。在CUDA当中,核函数kernel的执行总是异步的,而cudaMemcpy数据传输总是同步的。特别需要注意的是主机在提交核函数之后,不会阻塞等待核原创 2018-01-05 21:00:58 · 5977 阅读 · 0 评论 -
CUDA 流-重叠计算与数据传输
主机端启动kernel函数,可实现线程块级并行。为了实现任务级(线程格)级的并行,即同时在设备端启动多个核函数或数据传输操作,则需要引入CUDA流。一系列的设备任务(计算或数据传输)正是有序地排列在CUDA流中依次执行的,默认情况下所有任务都是执行在一个默认流中。当创建多个流时,可实现多个流的同步执行,以实现不同流中的各种操作达到重叠的效果,提升计算性能。cudaStream_t st原创 2018-01-05 21:20:27 · 1374 阅读 · 0 评论 -
CUDA-cudaEvent记录事件
CUDA中Event用于在流的执行中添加标记点,用于检查正在执行的流是否到达给定点。作用一,Event可用于等待和测试时间插入点前的操作,作用和streamSynchronize类似。作用二,Event可插入不同的流中,用于流之间的操作。不同流执行是并行的,特殊情况下,需要同步操作。同样,也可以在主机端操控设备端执行情况。作用三,可以用于统计时间,在需要测量的函数前后插入Event。调原创 2018-01-06 21:14:17 · 7266 阅读 · 0 评论 -
CUDA-非空流中的阻塞流
一.主机上非空流是异步流,其上所有的操作都不会阻塞主机执行。相应地,隐式的空流是同步流,大多数添加到空流上的操作都会导致主机在先前所有的操作产生阻塞。二.虽然非空流上在主机上是非阻塞的,但非空流内的操作可以被空流中操作所阻塞。因此可将非空流分为:阻塞和非阻塞两种。 如果非空流是阻塞流,则空流可以阻塞该非空流中的操作。 如果非空流是非阻塞流,则它不会阻塞空流中的操作。三,原创 2018-01-07 15:16:32 · 3701 阅读 · 5 评论 -
CUDA-同步
主机与设备之间的同步,分为隐式和显式。一.隐式:cudaMemcpy函数的作用在于传输传输,但在执行结束之前会产生阻塞。许多与内存相关的操作都会产生阻塞,这些不必要的阻塞会对性能产生较大的影响。如:锁页主机内存分配,设备内存分配,设备内存初始化,同一设备间的内存复制,一级缓存和共享存储配置的修改等等。二.显式:下面三种函数均可实现主机与设备间的同步。cudaDeviceSynchron原创 2018-01-07 16:09:12 · 1524 阅读 · 0 评论 -
CUDA全局内存-对齐与合并
CUDA执行模型的特征之一指令都是以线程束为单位进行发布和执行,存储操作也是如此。通常情况下,如图所示,全局内存是通过缓存进行加载,而加载请求是通过所谓“内存事务”来实现的,“内存事务”分为32字节和128字节两种。要么一次性加载32字节,要么一次性加载128字节,这主要取决于设备的缓存路径。如果对全局内存的访问使用L1/L2缓存存储(路径1),那么访问必须由128字节的“内存事务”完成,因此一行...原创 2018-01-11 20:29:08 · 5224 阅读 · 1 评论 -
CUDA全局内存读取
正如前文所述,CUDA全局内存的访问是通过”内存事务“实现的,其分类128字节(L1/L2缓存均参与)和32字节(L2缓存参与)两种。本文则主要介绍全局读取的加载示例,分为”缓存加载(L1+L2)“和”非缓存加载(L2)“,代码会贴在后面。一.缓存加载(L1+L2)这种情况下,”内存事务“中加载的缓存粒度是128字节。(1)对齐合并访问,线程束首地址对齐128字节,且连续访问128字节原创 2018-01-11 21:14:47 · 4234 阅读 · 0 评论 -
CUDA-全局内存读取-实验(缓存+非缓存-Fermi架构-sm2.1)
特别声明: 设备GT540M, 计算能力2.1.代码附在后面;缓存加载:(1)Fermi架构,默认情况是启用L1缓存,即采用128字节内存事务。 采用不同的偏移量,以实现非对齐访问。命令行为:“nvprof --metircs gld_efficiency test.exe N” (N为偏移量)。采用批处理,计算0-255的偏移量的全局内存加载效率,统计结果如下:偏移量每隔32,跳变原创 2018-01-11 21:18:18 · 1143 阅读 · 0 评论 -
CUDA-全局内存读取-实验(缓存+非缓存-Pascal架构-sm6.1)
特别声明: 设备GTX1050Ti, 计算能力6.1.代码附在后面;缓存加载:(1)Pascal架构,启用L1缓存,-Xptxas -dlcm=ca 。即采用128字节内存事务。 采用不同的偏移量,以实现非对齐访问。命令行为:“nvprof --metircs gld_efficiency test.exe N” (N为偏移量)。采用批处理,计算0-255的偏原创 2018-01-12 22:13:51 · 949 阅读 · 1 评论 -
CUFFT库(cufft_C2C,cufft_R2C,cufft_C2R,cufft_Z2C,cufft_D2Z,cufft_Z2D)
CUDA的cufft库可以实现(复数C-复数C),(实数R-复数C)和(复数C-实数R)的单精度,双精度福利变换。其变换前后的输入,输出数据的长度如图所示。在C2R和R2C模式中,根据埃尔米特对称性(Hermitian symmetry),变换后,*代表共轭复数。CUFFT的傅里叶变换类型则利用了这些冗余,将计算量降到最低。注意:下表都是单精度(C-表示float复数,R表示float实数)。而双原创 2018-01-18 00:23:24 · 7170 阅读 · 3 评论