
cuda&VS&matlab&arrayfire&GPU
文章平均质量分 52
ZIV555
GPU水声信号处理,电路设计
展开
-
GPU地址空间的相关概念
转载于:http://blog.youkuaiyun.com/xiewen_bupt/article/details/47208903知识点回顾:虚拟地址一个进程要对内存的某存储单元进行读写操作时,需要该内存单元的内存位置,这个时候,通过物理寻址计算出该内存位置。然而不同的进程可能修改相同位置的内存导致程序间的相互影响,因此,计算机系统引入了虚拟地址。每个进程都有自己的虚拟地址转载 2016-09-16 17:20:10 · 4442 阅读 · 0 评论 -
cuda debug 查看显存中变量
用工程自带的简单数组加法测试。一、打开start CUDA debug二、打开CUDA warp watch 及CUDA info三、就可以查看kernel中(显存)的变量了原创 2016-07-12 20:17:02 · 5222 阅读 · 1 评论 -
cuda及常用计时方式
CUDA中的计时方式:cudaEventCreate(&start1); cudaEventCreate(&stop1); cudaEventRecord(start1, NULL); Kernel>>( a, b, c ); cudaEventRecord(stop1, NULL); cudaEventSynchronize(stop1); cudaEvent原创 2016-06-22 15:23:53 · 1498 阅读 · 0 评论 -
cuda中warp分支及执行
根据编程指南,任何流控制指令(if\switch\do\for\while)都会导致线程分歧,线程分歧时会顺序执行每个分子路径,而禁用不在此路径上的线程,直到所有路径完成,线程再重新汇合到同一执行路径。主要有以下三种情况:如果一个warp的32个线程都只满足同一个分支,这是所谓的“分支按warp对齐”的形式,此时分支基本不影响执行效率。如果一个warp的32个线程分别满足多个分支,那么原创 2016-06-28 12:03:23 · 4214 阅读 · 0 评论 -
不同显卡(GPU)浮点数计算能力
1、SP总数=TPC&GPC数量*每个TPC中SM数量*每个SM中的SP数量;TPC和GPC是介于整个GPU和流处理器簇之间的硬件单元,用于执行CUDA计算。特斯拉架构硬件将SM组合成TPC(纹理处理集群),其中,TPC包含有纹理硬件支持(特别包含一个纹理缓存)和2个或3个SM,后面会有详细描述。费米架构硬件组则将SM组合为GPC(图形处理器集群),其中,每个GPU包含有一个光栅单元和4个SM原创 2016-06-24 16:59:15 · 26447 阅读 · 1 评论 -
每个线程块中最大线程问题
1:每个block 最大1024个线程(视不同的卡来定),这个是线程总数的限制。2:每个线程块最大维度的限制为x方向1024,y方向1024,z方向64(视不同的卡来定)。3:一个线程块的线程情况同时收到上述两条的约束,即,如果在x方向排布了1024个线程,那么y和z方向的维度只能是1,否则将超出第一条的约束。另外,每个SM的线程数也有最大规定,这个原创 2016-06-27 16:57:00 · 4255 阅读 · 3 评论 -
CULA安装
逛论坛 发现了这个东西留着备用 。CULA稠密矩阵 windows 版本:http://pan.baidu.com/s/1ntLTVq5CULA稀疏矩阵 windows版本:http://pan.baidu.com/s/1mg2tivY转载于http://bbs.gpuworld.cn/forum.php?mod=viewthread&tid=960转载 2016-06-06 09:30:44 · 1649 阅读 · 1 评论 -
利用matlab和NVIDIA Nsight进行cufft CUDA代码分析
分析之前肯定得有需要分析的代码,这里以fft为例:1、fft的C-mex cuda代码如下所示:#include "mex.h"#include #include void mexFunction (int nlhs,mxArray *plhs[],int nrhs,mxArray *prhs[]){ if(nrhs != 1) mexE原创 2016-06-03 17:16:53 · 2366 阅读 · 0 评论 -
独显和集显设置,双显卡工作及cuda计算
有时候我们会碰到插上独显后,想用集显却不显示了,不是电脑出问题了,需要在BIOS中进行设置。1、确保在BIOS中设置集显和独显都予以使用(部分板卡可能不支持);2、此时显示器可以接在集显上,独显单独做计算(确保都有合适的驱动);3、多显卡计算时,需要通过cudaSetDevice()来选定用来计算的GPU显卡。原创 2016-06-24 18:41:01 · 5347 阅读 · 1 评论 -
cuda stream处理
需要通过某种方式一次性地执行完读取、修改写入这三个操作,并且执行过程中不被其他线程中断,这种操作称为原子操作。#include "cuda_runtime.h"#include "device_launch_parameters.h"#include #include #include #define SIZE (100*1024*1024)__global__ voi转载 2016-06-24 17:46:47 · 1202 阅读 · 0 评论 -
visual profile 设置及用Visual Profiler进行CUFFT时间分析
上一篇博客中用nsight对cufft进行了kernel分析,这一篇仍然以cufft为例采用visual profile进行时间分析。1、现在你肯定已经完成的代码的撰写,cufft的代码和上篇博客一样,这里就不在赘述了。2、找到visual profile分析工具,如下图所示,双击打开:3、页面左上角点击file,然后file的子菜单中的new session,选择它,出现如下原创 2016-06-03 17:48:57 · 3358 阅读 · 0 评论 -
arrayfire中array转数组的问题
在arrayfire的葵花宝典中我暂时还没有发现直接转化的语句,只能先这样间接转化呢,仅供参考。 array in = randu(5, 5); af_print(in);for(int i=0;ifor(int j=0;j{vecl.push_back(sum(in(i,j)));}相当于对每个元素求和,还是自己。原创 2016-06-18 09:57:17 · 1289 阅读 · 1 评论 -
cuda中pinned memory(page-locked memory)
转载于:http://blog.youkuaiyun.com/zhangpinghao/article/details/21046435当为了提高CUDA程序的主机内存和设备内存传输消耗时,可以尝试一下两种方案一:使用分页锁定内存,分页锁定内存和显存之间的拷贝速度大约是6GB/s,普通的分页内存和GPU间的速度大约是3GB/s,(另外:GPU内存间速度是30G,CPU间内存速度是1转载 2016-08-04 10:36:33 · 1972 阅读 · 0 评论 -
cuda内存处理及stream内存处理
CUDA内存拷贝1、cudaMemcpy() cudaMalloc() //线性内存拷贝1 //线性内存拷贝2 cudaMalloc((void**)&dev_A, data_size);3 cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice);2、cudaMemcpy2D()cudaMallocPi转载 2016-07-09 13:31:16 · 1560 阅读 · 0 评论 -
GPU上基于SIMD的实现模式与多核CPU上基于MIMD的实现模式各有什么优缺点
如题: 由于每个执行单元的指令流都是相同的,SIMD模式将指令的获取时间均摊到每一个执行单元。但是,当指令流出现分支,指令就会被序列化。而MIMD模式的设计主要是为了处理不同指令流,当指令流出现分支,它不需要对线程进行阻塞。然而它需要更多指令存储以及译码单元,这就意味着硬件需要更多的硅,同时,为了维持多个单独的指令序列,它对指令带宽的需求也非常的高。 一般使用S原创 2016-06-23 16:31:58 · 3448 阅读 · 0 评论 -
Turbo Cache技术
Turbo Cache,即Turbo Cache技术,是NVIDIA推出的一项技术。通俗的说,Turbo Cache技术就是用内存当显存来使用。既然是利用系统内存,这和以前的集成显卡又有什么区别呢?在nVIDIA 的官方技术说明中,内存管理技术可以“允许GPU在分配和不分配系统内存时无缝切换,并且高效的读写内存”。这个工作由驱动程序中名叫TC的管理部分执行,以分配和平衡系统及本转载 2016-07-28 21:40:23 · 1294 阅读 · 0 评论 -
CBF中for循环变矩阵乘法的思想(arrayfire)--复数矩阵
接着上一篇博客CBF中for循环变矩阵乘法的思想(arrayfire)的续。上一篇主要讲了算法思想的改变,但是只是测试了实数,没有测试复数的效果,实际项目中都是复数的运用,所以这次添上复数的代码及测试结果。这次在添加arrayfire的代码之前,先看看不用这个库的一个C++代码形式:for(i=0;i<360;i++) //角度搜索 (-90:0.5:89) { theta=th原创 2016-06-22 13:33:44 · 1279 阅读 · 0 评论 -
安装CUDA调试器要关掉TDR
windows平台下配置parallel Nsight的第一步是禁用TDR功能。超时检测和恢复(TDR)是windows系统的一种机制,用来检测底层驱动代码的异常崩溃情形。如果驱动程序停止响应,windows则会重置此驱动程序。鉴于在程序断点处,驱动程序将暂停响应,所以为了防止出现重置操作或者GPU运行时间过长驱动自动恢复(修改TDR中已经提到了),TDR功能需要关掉。怎么关掉如下图所示:原创 2016-06-21 10:57:38 · 3628 阅读 · 1 评论 -
CBF中for循环变矩阵乘法的思想(arrayfire)
/******************************************************* * Copyright (c) 2016.6.19 * 算法的改进:主要解决了两个for循环的问题 *第一个for是角度的问题,循环360个角度或者更小来得到每个角度的 *输出功率。 *第二个for循环是通道的问题,目前是8通道,对于arrarfire来讲这个很好做 *直接原创 2016-06-19 11:33:14 · 1305 阅读 · 0 评论 -
Arrayfire常用的那几招(引用于葵花宝典)
/**************************arrayfire在矩阵运算中常用的那几招*2016.6 ziv*************************/seq:线性序列,主要用来替换for循环eg:seq b(10, 20, 2); // [10, 20, 2] => 10, 12, 14....20其中10-begin,20-end,2-ste原创 2016-06-19 16:08:47 · 3962 阅读 · 1 评论 -
不同工具下的矩阵乘法速度测试
今天花了一些时间将基于以上几种工具(cuda&arrayfire&matlab)的矩阵乘法的速度进行了测试比较,验证了一些想法吧。首先是c(CPU)的乘法测试:写的有点繁琐,后面在cuda程序中进行了综合C.matrix.cppvoid Matrix_print(double **a,long nl, long nh)/*矩阵的输出*/{ int i,j; for(i=1;原创 2016-06-19 16:53:06 · 2275 阅读 · 0 评论 -
关于线程的问题
其实这个问题纠结很久了,很多版本也不一样,现在也不是很清楚,把现在的情况整理一下吧;1、首先通过通过GPU-Z或者CUDA-Z会看到有多少个流多处理器、多少个流处理器、每个流多处理器会有多少个threads,每个block至少多少个threads等,关键是线程的分配问题,是不是一个sp(流处理器)执行一个线程,那么多少个sp就只能执行多少个线程,不是这样的。2、首先分配上根据一个SM最大线原创 2016-06-03 11:07:01 · 633 阅读 · 0 评论 -
矩阵乘法中的高阶计算时间和计算误差问题
主要探讨的有两个问题:1、当矩阵维度为16*16,2048*2048时,可以计算出来(用全1矩阵进行测试),并且加速比大于1(GPU快),可是当矩阵维度提升到4096*4096时,此时会出错(计算不准确,显示器驱动程序已停止响应 并且已成功恢复 ),怎么解决?2、先前测试了全1矩阵,现在改成随机数矩阵(包括小数和整数),16*16等矩阵都不能得到正确的结果,怎么办?首先贴出多线程的原创 2016-05-31 14:03:03 · 3118 阅读 · 0 评论 -
关于 Nsight Monitor 启动时出现port8000 怎么解决
因为需要打开Nsight Monitor调整TDR,可是每次打开总是会粗线,下面情况。首先我们得找出占用这个端口的值。1、首先我们打开CMD,用管理员身份打开,然后输入netstat -ano,列出所有端口的情况。2、查看被占用端口对应的PID,输入命令:netstat -aon|findstr "8000",回车,记下最后一位数字,即PID,如果是2839。3、(可用可不用)原创 2016-05-31 14:33:06 · 1714 阅读 · 0 评论 -
False dependencies
留着以后备用是一种在N卡上,哪怕使用了多个stream, 却取不到并行的效果(或者能取到,但有限)。这在某些N卡和N卡架构,以及驱动版本配合上,常见的。您在具有计算能力3.5或者更高的Tesla显卡上,不应当能容易观察到这种现象的。(除非你限定了CUDA_DEVICE_MAX_CONNECTIONS)这个问题的产生是因为某些N卡的硬件,发现独立无关的并行任务(包括复制传转载 2016-04-15 10:36:24 · 820 阅读 · 0 评论 -
kernel的红色波浪线
有时候 kernel_4 > >()括号下面会出现红色波浪线,是VS没有识别,但是不妨碍你程序运行,不用管它。原创 2016-04-15 10:28:17 · 543 阅读 · 0 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(一)——序
看看这个题目,我自己都觉得比较混乱! 先说说与这个有关的方向,水声信号处理(并行波束形成),开始接触这方面应该是 去年10月份的时候吧,要做一个实时波束形成的东西,那时候并不知道GPU在信号处理的运用,在大哥和厂家的共同推荐下,就走上了这条路,现在看看,因为能力有限,好像并没有取得预期的效果(只是用arrayfire做出了大概的样子,但是和MFC的融合我又fuck了原创 2016-04-06 10:21:43 · 1962 阅读 · 3 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(二)——matlab
Matlab&GPU&CUDA并行加速学习心得1所需软件:l NVIDIA GPU一块或多块,允许不同型号的GPU共存(本机配置了华硕GTX960);l NVIDIA开发驱动程序(安装板卡的时候会自动提示安装,没有提示安装用附带光盘);l CUDA toolkit安装,这里选择的是cuda toolkit7.5安装,网址如下,里面包含了GPU计算软件开发工具包,并行Nsigh原创 2016-04-06 10:26:53 · 2184 阅读 · 0 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(三)——cuda
突然想找到这个文章的链接,但是可惜找不到了忘记了版主是谁呢,但是这篇文章比较系统的介绍了CUDA方面的知识。(忘记说了我使用的电脑配置gtx970 和 980都用过,toolkit7.5)1 CUDA C编程入门-介绍1.1 从图形处理到通用并行计算在实时、高清3D图形的巨大市场需求的驱动下,可编程的图形处理单元或者GPU发展成拥有巨大计算能力的和非常高的内存带转载 2016-04-06 10:42:16 · 2901 阅读 · 1 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(四)——matlab mex function
主要介绍matlab转c的功能,然后mex成.mexw64(我的电脑是64位的)matlab中有matlab coder这个工具,可以直接转换成.cpp文件 ,mex function 可以直接将其转化成.mexw64;需要注意的一点是必须将需要转换的算法模块都写成function,m文件的开头必须是function[输出参数1,输出参数2,...]=m文件名(输入原创 2016-04-06 21:50:37 · 906 阅读 · 0 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(五)——square的例子
这个地主要学习了《实战matlab并行程序设计》中的square这个例子先建一个CUDA工程,然后.cu文件,把代码敲入#include"mex.h"#include"cuda.h"#include "cuda_runtime.h"#include "cuda.h"#include "device_launch_parameters.h" #include原创 2016-04-06 21:52:08 · 1057 阅读 · 0 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(六)——c语言转换
主要是常规波束形成和方位历程图;常规波束形成的matlab部分代码如下所示:for i = 1:length(theta) a_s = exp(-jay*2*pi*[0:N-1]*f0*d/c*sin(theta(i)*pi/180)); (大哥写的) beam_single(i) =1/(a_s *inv(R1) *a_s');end 首先在网原创 2016-04-06 21:57:23 · 1462 阅读 · 0 评论 -
GPU&VS2012&CUDA&matlab&Arrayfire杂记(七)——MFC语言与CUDA
CUDA 与 MFC环境:vista,vs2005,cuda2.0前提工作:创建一个MFC工程,本例名称为cudaexam说明:本文件中的方法仅仅是向导之一,且没有进行并行优化工作。步骤:1, 创建全局函数与头文件1.1头文件1.2函数1.3然后查看类视图2, 创建一个存放cuda代码的筛选器,本程序名称为cuda2.12.2转载 2016-04-06 22:02:38 · 1051 阅读 · 0 评论 -
过大的数组该怎么存储
(1)开巨大的stack。(你至少需要7560000 * 8 大约 60MB的stack !!)(这种方法不推荐)(2)不使用巨大的stack, 而改用heap分配,例如通过malloc或者new之类的。(这是对于host code说的)转载于http://bbs.gpuworld.cn/forum.php?mod=viewthread&tid=10223&extra=pa转载 2016-04-15 10:50:02 · 3384 阅读 · 0 评论 -
开普勒架构和麦克斯韦架构是什么?有什么区别?
Maxwell可以算Kepler的改进版架构。两个架构最明显的变化是在SMX单元和GPC单元上。Maxwell的SMM(之前叫SMX)单元从之前Kepler的包含192个CUDA Core下降到128个,但发射器从之前的每SMX一个变为了每SMM四个,目的是降低每个SMM单元的运算压力提升效率。增加了两个寄存器,然后L1缓存翻倍,GPC单元的L2缓存增加到了2M。现在已经上市的Maxwel转载 2016-05-10 09:36:50 · 13302 阅读 · 0 评论 -
GPU(显卡)的WDDM TDR时间修改方法
为了暂时利用GPU(显卡)进行计算,需要对桌面显示作出一个延时,在矩阵乘法上遇到当计算时间过长时会遇到下面这种情况,这是计算时间过长的表现,此时就需修改TDR的时间,如是有了这篇转载的文章。1, 联通两台电脑:准备两台电脑,分别称为主机和调试机,在主机端运行vs2010,在调试机上进行调试,其中至少调试机应支持CUDA,使两台机器在同一个局域网,或直接将两台电脑用网线连接起来转载 2016-05-31 14:12:48 · 12216 阅读 · 0 评论 -
CUDA数组学习
在群里看到的,因为开始也遇到过类似问题,所以转载过来备份一份。1:问题是怎么来的在device上要用一维数组、二维数组、三维数组,对于一维数组用了cudaMalloc和cudaMemcpy进行内存分配和赋值,但是对于二维和三维的分配本想这样转换成一维进行。但是这样感觉赋值又不方便,刚刚看了一个例子,这样做的:[cpp] view plain copy转载 2016-04-22 16:13:16 · 1716 阅读 · 0 评论 -
常规波束形成的CPU和GPU的运行时间对比测试
如图所示,做了一点小测试,确实和以前看到的资料一样,会出现临界数据量,但是我觉得我的GPU程序写的不够好吧,没有达到想要的效果,不知道大家有什么看法。然后通过此次测试认为影响GPU加速的几个因素:1:程序中出现较多的循环会拖慢速度,我用了while,所以将这些东西转换成向量或许会快一些;2:数据在传输的时候耗时比较多,内存优化做的不够好;3:因为用的ar原创 2016-04-22 16:11:10 · 1546 阅读 · 0 评论 -
matlab中图线颜色大全
(matlab)plot画图的颜色线型 y 黄色 · 点线 m 粉红 ○ 圈线 c 亮蓝 × ×线 r 大红 +转载 2016-04-22 09:57:10 · 10292 阅读 · 0 评论 -
GPU中while问题
刚才上面那个博客忘记说了刚才测试的问题,在做方位历程图的时候,需要将数据分段,然后再对每段数据做方位谱估计。我现在的数据量为1024*64,当把数据分成103段的时候,cpu时间为4.9s,gpu时间为6.9s(可能我写的程序很挫);但是这里想强调的是,当分段为1047段后,cpu运行时间为8.45s,而GPU运行时间为68s; 不得不说这个差距太大了,造成这个极速变慢的原因可原创 2016-04-21 19:25:07 · 775 阅读 · 0 评论