Matlab下的CUDA编程(二)

2010年02月22日 by wy,59 views

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

/*
        This is a mex file to offload  the FFT2 function in Matlab to CUDA:
 
    Y = fft2(X)
    or
        Y = fft2(X,M,N)
 
    Y = fft2(X) returns the two-dimensional FFT of X. The result Y is the same size as X.
 
    Y = fft2(X,M,N) truncates X, or pads X with zeros to create an M-by-N array before doing the transform.
        The result is M-by-N.
 
        On Windows:
 
    Setup the mex build from a Matlab shell:
    mex -setup
 
        Compile the mex file:
    mex -IC:\CUDA\include fft2_cuda.c -LC:\cuda\lib -lcufft -lcudart
 
        On Linux:
        Use the makefile
 */
 
#include "mex.h"
#include "cufft.h"
#include "cuda_runtime.h"
 
/**************************************************************************/
 
/* MATLAB stores complex numbers in separate arrays for the real and
   imaginary parts.  The following functions take the data in
   this format and pack it into a complex work array, or
   unpack it, respectively.
   We are using cufftComplex defined in cufft.h to  handle complex on Windows and Linux
*/
 
void pack_r2c(cufftComplex  *output_float,
                double *input_re,
                int Ntot)
{
    int i;
    for (i = 0; i < Ntot; i++)
    {
        output_float[i].x = input_re[i];
        output_float[i].y = 0.0f;
    }
}
 
void pack_c2c(cufftComplex  *output_float,
                double *input_re,
                double *input_im,
                int Ntot)
{
    int i;
    for (i = 0; i < Ntot; i++)
    {
        output_float[i].x = input_re[i];
        output_float[i].y = input_im[i];
    }
}
 
 
void pack_r2c_expand(cufftComplex  *output_float,
                double *input_re,
                int originalM,
                int originalN,
                int M,
                int N)
{
    int i, j;
    for (i = 0; i < originalM; i++)
        for (j = 0; j < originalN; j++)
    {
        output_float[i+M*j].x = input_re[i+originalM*j];
        output_float[i+M*j].y = 0.0f;
    }
}
 
void pack_c2c_expand(cufftComplex  *output_float,
                double *input_re,
                double *input_im,
                int originalM,
                int originalN,
                int M,
                int N)
{
    int i, j;
 
    for (i = 0; i < originalM; i++)
        for (j = 0; j < originalN; j++)
    {
        output_float[i+M*j].x = input_re[i+originalM*j];
        output_float[i+M*j].y = input_im[i+originalM*j];
    }
}
 
void unpack_c2c(cufftComplex  *input_float,
                double *output_re,
                double *output_im,
                int Ntot)
{
    int i;
    for (i = 0; i < Ntot; i++)
    {
        output_re[i] = input_float[i].x;
        output_im[i] = input_float[i].y;
    }
 
}
 
 
/**************************************************************************/
 
void mexFunction( int nlhs, mxArray *plhs[],
                int nrhs, const mxArray *prhs[])
{
    int originalM, originalN, M, N;
    double *ar,*ai;
    cufftComplex *input_single ;
    cufftHandle            plan;
    cufftComplex *rhs_complex_d;
 
    /*
     Find out the  dimension of the array we want to transform:
 
     prhs(originalM,originalN)
     originalM = Number of rows    in the mxArray prhs
     originalN = Number of columns in the mxArray prhs
    */
 
    originalM = mxGetM(prhs[0]);
    originalN = mxGetN(prhs[0]);
 
    M = originalM;
    N = originalN;
 
    /*
    Find out if the result is of the same size of the input.
 
    plhs(M,N)
    M = The number of rows in the mxArray plhs
    N = The number of columns in the mxArray plhs
    */
    if (nrhs == 3)
    {
        M = mxGetScalar(prhs[1]);
        N = mxGetScalar(prhs[2]);
    }
  /*
    Find out if the input array was real or complex.
    Matlab is passing two separate pointers for the real and imaginary part.
    The current version of  CUDAFFT is expecting interleaved data.
    We will need to pack and unpack the data.
 
    The current version of CUDA supports only single precision,
    so the original double precision data need to be converted.
   */
 
  /* Allocating working array on host */
    if(nrhs == 1) input_single  = (cufftComplex*) mxMalloc(sizeof(cufftComplex)*N*M);
    if(nrhs == 3) input_single  = (cufftComplex*) mxCalloc(N*M,sizeof(cufftComplex));
 
  /* Pointer for the real part of the input */
    ar =  (double *) mxGetData(prhs[0]);
 
    if (mxIsComplex(prhs[0]))
    {
   /* The input array is complex */
        ai =  (double *) mxGetImagData(prhs[0]);
 
   /* Input and output have same dimensions */
        if(nrhs ==1) pack_c2c(input_single, ar, ai, N*M);
 
   /* Input and output have different dimensions */
        if(nrhs ==3) pack_c2c_expand(input_single, ar, ai, originalM, originalN, M, N);
 
    }
    else
    {
   /* The input array is real */
 
   /* Input and output have same dimensions */
        if(nrhs ==1) pack_r2c(input_single, ar, N*M);
 
   /* Input and output have different dimensions */
        if(nrhs ==3) pack_r2c_expand(input_single, ar, originalM, originalN, M, N);
    }
 
  /* Allocate array on device */
    cudaMalloc( (void **) &rhs_complex_d,sizeof(cufftComplex)*N*M);
 
  /* Copy input array in interleaved format to the device */
    cudaMemcpy( rhs_complex_d, input_single, sizeof(cufftComplex)*N*M, cudaMemcpyHostToDevice);
 
  /* Create plan for CUDA FFT */
    cufftPlan2d(&plan, N, M, CUFFT_C2C) ;
 
  /* Execute FFT on device */
    cufftExecC2C(plan, rhs_complex_d, rhs_complex_d, CUFFT_FORWARD) ;
 
  /* Destroy plan */
    cufftDestroy(plan);
 
  /* Copy result back to host */
    cudaMemcpy( input_single, rhs_complex_d, sizeof(cufftComplex)*N*M, cudaMemcpyDeviceToHost);
 
    plhs[0]=mxCreateDoubleMatrix(M,N,mxCOMPLEX);
 
    ar = mxGetPr(plhs[0]);
    ai = mxGetPi(plhs[0]);
    unpack_c2c(input_single, ar, ai, N*M);
 
    mxFree(input_single);
    cudaFree(rhs_complex_d);
 
    return;
}

Matlab下的CUDA编程(一)

2010年02月4日 by wy,111 views

引言

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

AMD/ATI 发布Stream SDK 2.0正式版

2009年12月23日 by annkok,104 views

详细报道:
Link1
Link 2
官方下载
大致看了以下,依旧没有提供对image的支持,不过好歹增加了对dx和opengl的互操作。
根据之前测试版的经验,对性能提升不太报希望(ps 之前测试性能竟然和brook+实现在一个级别,比较汗的说)。

基于CT重建的性能测试稍候提供。

Larrabee杯具了

2009年12月6日 by annkok,192 views

Intel说larrabee要以软件平台的形式发布了,至于实体…等吧。
Nvidia和ATI又可以小爽一段了。

话说现在竞争白热,局势却相对明朗了:Nvidia狂推自己的CUDA,AMD和IBM等一干人紧紧抱着Opencl,Intel自己没搞出来东西,又不甘落后,就自己再搞一套软件平台来搅浑水,Microsoft肯定是不屑跟他们这些人玩了,肯定要抱着自己的DX Compute闷头折腾了。

我们这些下游就得屁颠跟着他们晕了。
新闻Link

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

2009年12月6日 by emaxin,186 views

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

OpenCL Benchmark

2009年12月5日 by annkok,139 views

AMD和Sisoft联合在Sisoft的Sandra 2010里面添加了opencl benchmark。支持的硬件有:

  • ATI GPUs: Radeon HD 5800, 5700 series
  • AMD CPUs: Athlon Neo, Six (6)/Eight (8)/Twelve (12)-Core Opteron 4100, 6100 server
  • Intel CPUs: Core i5/i7 desktop/mobile, Xeon Nehalem-EX server
  • nVidia GPUs: GT215, GT216, GT218 series
  • VIA CPUs: Nano 3000 series
  • Hardware monitoring: Analog Devices (ADT), SMSC, Winbond/Nuovoton

大致是对带宽的测试和计算性能的测试,结果:
GPU结果
CPU结果
比较诡异的是,amd的显卡是5870 ,NVidia的显卡竟然选用的是9500,这个…有点太…

OpenCL系列(四)-Memory

2009年11月28日 by annkok,149 views

Memory分类和访问(Memory Layout and Access)

opencl将存储模型分为三种:Private Memory,Local Memory,Constant Memory和Global Memory。其中Private Memory只可以被work-item使用,也就是单个work-item私有的;Local Memory可以被同一个work-group内的work-item所使用;Global Memory可以被所有的work-item访问和修改;而Constant Memory是Global Memory的一种,区别在于它是不能被更改的。示意图如下[1]:
opencl_memory_layout

大致了解了opencl中存储类型后,就会发现opencl的memory model和cuda的极为类似,ok,这里就简单的列出opencl中集中memory类型在Nvidia的GPU、AMD的GPU和CPU中实际对应的数据类型:

Opencl Nvidia CUDA ATI GPU AMD CPU
Global Memory Global Memory global buffer (g[] in CAL/IL) 系统内存(RAM)
Local Memory Shared Memory Local Data Share CPU L1缓存
Private Memory Register scratch memory (x[] in CAL/IL) 系统内存RAM

update@2009.11.30
Opencl中还有一种存储类型:Image,也就是cuda中的纹理内存,包含2D和3D image。与CUDA中纹理一样提供不同的插值(采样)方式,同时也提供一定cache(硬件支持的),个人感觉,这是一个GPU的遗留产物,在CPU和DSP等的Opencl实现极有可能不提供对其的支持。所以要注意在使用时,最好先查询一下是否支持。

如何使用(How to use these memory)
to be continued…

IBM停止CELL开发

2009年11月22日 by annkok,262 views

坑,又少了一个。原文
update:好像是说停止了相关sdk的开发,转向了opencl,这点和前几天从amd那边得到的消息一致。

Ubuntu 9.04 安装 CUDA 开发环境

2009年11月13日 by pt,317 views

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) 阅读这个条目剩下部分 »

OpenCL系列讲座(三) Step By Step

2009年11月8日 by emaxin,287 views

  我们已经写好的Kernel函数一般保存在一个字符串内,例如在第二节中的kernel:输入两个一维数组,用另外一个一维数组输出两个一维数组求和的结果,如下:

char* chKernelSource = "__kernel sum 
(__global const float* a, __global const float* b,  
__global float* answer)
{
   int xid = get_global_id(0);
    answer[xid] = a[xid] + b[xid];
}";

概述中所言,下面我们一步一步来写Host代码Main函数部分;
  阅读这个条目剩下部分 »