文章关键字 ‘CUDA’

cuda Array 拷贝问题

2010年05月3日,星期一

今天作者在使用cudaMemcpyArrayToArray函数时发现其效率极低,如果使用如下代码,需要77ms

cudaMemcpyArrayToArray(d_Array_A, 0, 0, d_Array_B, 0, 0, nFrameSize*sizeof(float));

如果开辟一块显存中转一下,效率会大幅提高,到20ms

cudaMemcpyFromArray(d_pfData, d_Array_B, 0,0,nFrameSize*sizeof(float),cudaMemcpyDeviceToDevice);
cudaMemcpyToArray(d_Array_A, 0, 0, d_pfData, nFrameSize*sizeof(float), cudaMemcpyDeviceToDevice);

纹理内存与全局内存之间有缓存,所以数据交换较快可以理解,但纹理内存之间拷贝速度反而更慢则不能理解。估计是cuda内部实现时没有做好优化。

Matlab下的CUDA编程(四)

2010年04月7日,星期三

第三节中我们介绍了NVIDIA的工程师写的编译脚本,但是这个脚本配置起来麻烦,用的时候选项也过于复杂。后来随着GPGPU的发展,CUDA逐渐被重视,因此MathWorks的工程师们重新写了这个nvmex脚本。毕竟是科班出身,新的脚本配置容易,使用简单。详细情况可访问其网站。在上述网站中,可找到下载nvmex源码的链接。下载后解压,得到nvmex.m。对其中的两个选项(红色部分按实际路径修改)稍加编辑,即可使用:

CUDA_LIB_Location = ‘C:\CUDA\lib’;
Host_Compiler_Location = ‘-ccbin "C:\Program Files\Microsoft Visual Studio 8\VC\bin"‘;

修改之后,将其复制到addMatrix.cu相同目录,并将此目录设置为matlab运行目录,在命令窗口输入:

>> nvmex(‘addMatrix.cu’);

即可完成编译。编译成功后,在matlab中即可像常规函数一样使用addMatrix函数。

Matlab下的CUDA编程(三)

2010年04月7日,星期三

根据第二节的介绍,matlab中可以通过mex文件的方式编译C/C++代码,但是对于.cu文件则无能为力。为了解决这个问题,NV的工程师们开发了用于编译cu文件的脚本,下面进行详细介绍。 测试环境:vs2005, matlab 7.6(r2008a), cuda 2.3 首先在NV的网站http://developer.nvidia.com/object/matlab_cuda.html下载matlab的插件包,并解压,共有4个重要文件:

nvmex.m
nvmex_helper.m
nvmexopts.bat
.\bin\nvmex.pl

首先将nvmex.pl拷贝到matlab的安装目录中的bin目录下,例如笔者的系统中就拷贝到C:\MATLAB\bin中。 下面编写cu文件,我们仍以两个矩阵相加为例,具体流程见代码注释。从代码中可看到,同时使用了matlab函数和cuda函数,并在GPU中完成了两个矩阵相加。完成addMatrix.cu的编写后,需要对其进行编译。 首先将nvmex.m,nvmex_helper.m,nvmexopts.bat拷贝到addMatrix.cu文件所在目录,并设置为matlab的当前目录。

  1. 在matlab命令窗口中运行: nvmex -setup,选择编译器,此处选择vs2005。
  2. 编辑nvmexopts.bat,修改其中的VSINSTALLDIR选项,指定为vs2005的安装目录,如笔者的系统中: C:\Program Files\Microsoft Visual Studio 8
  3. 编译命令: nvmex -f nvmexopts.bat addMatrix.cu -IC:\cuda\include -LC:\cuda\lib –lcudart

编译成功后,在matlab中即可像常规函数一样使用addMatrix函数。 可能会遇到的问题:

  1. 目录设置:需要仔细按照上述步骤设置好相关目录
  2. 编译时出现未定义的变量及结构,如_wchar等,则可更新nvmex.pl和nvmexopts.bat来解决。下载链接

(全文…)

CUDA 3.0正式发布

2010年03月20日,星期六

CUDA 3.0正式发布。下载链接。一些新的features(来自release notes,按照惯例稍微评论一下):

    新的API:

  • Half float(float16)纹理支持
    对一些精度要求不高的算法,可以copy更多的数据到显卡了,比如CT重见算法,已经有人写paper证明了:)
  • double3 double4 矢量类型支持
    没啥可说的了,就新的支持
  • 一维的设备到设备的拷贝支持streams了
    (确实没有感觉到stream的好处,之前折腾过一段,效果不明显)
  • 二进制ELF支持
  • Concurrent Kernels(通过CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS 这个属性值才查看)
    貌似这个比较关键,也比较好,只是实际用起来不知道效果如何?
  • 批量的2D 3D FFT支持
    不做评论,暂时没有用到
    新的Toolkit功能增加

  • Nvcc的–host-compilation=C选项没有了
    这个没啥说的,难道这就是Nvidia说的原生的C++支持???
  • Window下cuda的DLL命名规定,也就是规范了下cudart.dll的名字,允许使用多个cudart了。比如cudart32_30_4.dll 就是32位3.0版的dll了,4是编译的个数??唉不会想那些vc运行库那样把
  • 将模拟器模式从cudart.dll单独分离出来了,以后用模拟器就要用cudaemu了
  • CUBLAS 新的函数
    没用到,不关注

(全文…)

Matlab下的CUDA编程(二)

2010年02月22日,星期一

Matlab下通过Mex文件编写C程序
本节参考NVIDIA网站相关资源,点击此处链接:

1、MEX规则

Matlab提供了MEX文件的方式来支持C/C++代码编写的算法。mex文件需要满足如下要求:
(1) 包含mex.h头文件
(2) 函数名称、参数返回值必须为如下形式:

void mexFunction(int nlhs, mxArray *plhs[],int nrhs, const mxArray *prhs[]);

其中:
nlhs 为输出数组个数(Left Hand Side)
plhs 为指向输出数组的指针
nrhs 为输入数组个数(Right Hand Side)
prhs 为指向输入数组的指针,且输入数组只读。
以上四个变量,是在C/C++代码中唯一可用到的变量。实际上,由于matlab中所有的变量都是mxArray结构(向量,数组,字符串。。。),因此常见的数据类型均可放入mxArray传递给程序进行处理。
(3)常用mex函数:
a. mex函数定义在mex.h(./extern/include)中,并以mex为前缀,例如打印输出的mexPrintf()函数、在mex文件中调用matlab函数的mexCallMATLAB()函数等。
b. 在matrix.h(./extern/include)中定义了mxArray结构以及对矩阵操作的函数,如创建双精度矩阵的mxCreateDoubleMatrix()、从mxArray中获取数据指针mxGetPr()等。
(4)编译
首先可以通过mex -setup来选择编译器,此时matlab会提示:
Would you like mex to locate installed compilers [y]/n?
此处如果选y,则会列出系统中安装的编译器,但不一定完整(作者就遇到这样的情况)。如果选n,则会列出matlab支持的所有编译器,
我们按实际情况选取即可(按照CUDA的要求,vs2005以上)。
另外常用的编译选项有:
-I 增加头文件(.h)包含目录
-L 增加库文件(.lib)包含目录
-l 链接引用的库文件名称

其他命令可查看matlab中的帮助文档。

2、MEX例程
在本教程中,举例实现将两个矩阵相加(C = A + B)。

/********************************************************************
filename:     addMatrix.c
file ext:       c
author:        wy@gpgpu.org.cn
purpose:  test matlab with c
*********************************************************************/
#include "mex.h"
void mexFunction(int nlhs, mxArray *plhs[],int nrhs, const mxArray *prhs[])
{
    int i, j, mA, nA, mB, nB, nMatSize;
    double *A, *B, *C;
    /* 输入变量检查 */
    if (nrhs != 2)
        mexErrMsgTxt("输入参数必须为2个");
    if (nlhs != 1)
        mexErrMsgTxt("输出参数必须为1个");
    mA = mxGetM(prhs[0]);
    nA = mxGetN(prhs[0]);
    mB = mxGetM(prhs[1]);
    nB = mxGetN(prhs[1]);
   if (mA != mB ||
       nA != nB)
       mexErrMsgTxt("输入矩阵尺寸必须相同");
    /*为输出结果分配内存*/
    plhs[0]=mxCreateDoubleMatrix(mA,nA,mxREAL);
    /*获取数据指针*/
    A = mxGetPr(prhs[0]);
    B = mxGetPr(prhs[1]);
    C = mxGetPr(plhs[0]);
    /*求和计算*/
    nMatSize = mA * nA;
    for (i=0; i<nmatsize ; i++)
    {
        C[i] = A[i] + B[i];
    }
    /*输出信息*/
    mexPrintf("求和完毕");
}

编译c文件:
mex -v addMatrix.c

编译成功后,在matlab中即可像常规函数一样使用addMatrix函数。
a = rand(2,2)
a =
0.93547      0.41027
0.9169      0.89365

b = rand(2,2)
b =
0.057891      0.81317
0.35287    0.0098613

c = addMatrix(a,b)
求和完毕
c =
0.99336       1.2234
1.2698      0.90351

3. 使用cuda标准库
在nvidia给出的白皮书例程中,给出了二维FFT的实现方法,本文将其引用,说明通过mex文件使用cuda标准库实现加速的方法。
在该程序中,没有调用内核程序,而只使用了runtime库和fft库,并不需要.cu文件,可以用mex命令直接编译。因此,将cuda算法封装之后再由mex文件调用是使用matlab使用cuda的一种方法。
编译命令:
mex -IC:\CUDA\include fft2_cuda.c -LC:\cuda\lib -lcufft -lcudart
(全文…)

Matlab下的CUDA编程(一)

2010年02月4日,星期四

引言

Matlab作为科学计算中的重要工具,它提供了丰富的基本函数和工具箱,从而在各个领域得到了广泛的应用。但Matlab的缺点是代码效率较低,在工程应用中只能作为模型验证,而不能得到实际应用。而CUDA作为显卡编程比较成熟的语言,能够充分利用GPU的计算能力,提高执行效率。因此如何能将CUDA的高效与matlab的简便有机的结合是本文要解决的问题,根据解决问题的方式我们将分为两节进行讲解。

一种简单的GPU三维图像分割算法

2009年12月6日,星期日

图像分割的定义是将整幅图像区域分割成各自互不交叉的小区域。图像分割的算法有很多,教科书上讲到一些经典的分割算法,包括阈值法、区域生长法、分裂与合并算法、边界跟踪与连接算法、分水岭算法,这些算法适合提取出灰度比较一致的各个区域;此外还有纹理分割算法,严格说来,与其称之为纹理分割还不如称之为纹理检测与分类,通过提取纹理特征,如:Laws特征、梯度共生矩阵特征、分形特征,得到每一个点或小邻域的特征向量,然后利用特征向量将每个点分到不同的纹理类中去。
有些扯远了,还是回到三维图像分割上来,由于三维图像的数据量是巨大的,暂不考虑复杂的纹理特征,我们仅仅希望有在GPU上实现一种并行的快速算法,能够自动标记出三维数据中所有灰度比较一致的各个连通区域。我们先看看教科书上的方法,阈值法不能提取多个灰度不一致的区域,区域生长不容易做到每个区域种子点自动选取,分裂与合并的逻辑稍显复杂且分割的区域边缘会出现锯齿,边界跟踪与连接在三维上根本没办法实现,分水岭又是串行递归的-_-!!!
假设已经是有了一幅二值图,我们就可以用连通区标记的算法,为二值图中的每一个小的区域标上不同的标号,这就是常用的二值图CCL(Connected Component Labeling)算法,可以扫描填充法、FloodFill等方式来具体实现这个算法。算法的原理是,判断当前点的邻域内有没有和当前点取值相同的点,如果有,就把他们标记为相同的标号。利用这个思想,如果在判断的时候,不是判断当前点的邻域内有没有和当前点取值相同的点,而是判断当前点的邻域内有没有和当前点取值相近的点——有点类似于梯度区域生长算法,就可以设计出适用于灰度图像的灰度连通区标记算法灰度图CCL算法,如果对每一个点进行如上判断,就可以预先生成一幅具有二值图,且对每个点的判断不依赖于其它点判断的结果,可以做到完全并行。这一段代码很简单,假设体数据已经做了去噪处理,然后用For循环遍历每一个点就能搞定(对于Cuda实现,就是把最里层的内容写到kernel里面),贴代码:
(全文…)

Ubuntu 9.04 安装 CUDA 开发环境

2009年11月13日,星期五

1. 下载 CUDA 文件

从 Nvidia CUDA 官方网站( http://www.nvidia.com/object/cuda_home.html )的下载页面可以找到适用于 ubuntu 9.04 的 CUDA 文件,包括:

*  NVIDIA Driver 190.18 Beta for Linux (Ubuntu 9.04) with CUDA Support
*  CUDA Toolkit 2.3 for Linux (Ubuntu 9.04)
*  CUDA SDK 2.3 code samples for Linux (Ubuntu 9.04) (全文…)

CUDA 3.0

2009年11月6日,星期五

The CUDA Toolkit 3.0 Beta is now available to GPU Computing registered
developers

OpenCL系列讲座(一) 概览

2009年11月4日,星期三

OpenCL(Open Computing Language)是由Apple发起,由Khronos负责起草的用于跨平台并行计算的工业标准。基于该标准,我们可以使用任何系统中支持该标准的计算资源。由于该标准支持不同的编程语言、不同的硬件设备并且高效,因此可能成为软件、中间件乃至使用高性能计算设备厂商的有效解决方案。
(全文…)