matlab中使用CUDA kernel GPU加速

本帖最后由 蓝云风翼 于 2013-6-8 14:13 编辑

截至 MATLAB2013a 里面已经有不少工具箱里面都有了支持GPU加速的函数。使用matlab+GPU加速的前提是机器必须安装了支持CUDA的显卡，且GPU 计算能力在1.3以上。
支持的GPU 可通过gpuDevice 查看GPU是否支持
支持GPU加速的函数可通过methods(‘gpuArray’)查看

例如fft，ifft，三角函数，相关函数xcorr以及常用的运算符等等都可以进行加速。方法也很简单，主要使用到gpuArray和gather这两个函数。
以xcorr为例，假设我们要求向量A和B的互相关，一般是使用代码
M = xcorr(A,B)
以下是使用gpu加速的版本
Ag = gpuArray(A);
Bg = gpuArray(B);
Mg = xcorr(Ag,Bg);
M = gather(Mg);

当然本帖不主要介绍如何使用gpuArray，本帖将介绍如何在matlab编译.CU文件生成PTX文件，对于nvcc，ptx编译原理有兴趣的可参看 CUDA编程接口：如何用nvcc编译CUDA程序
内核可以使用PTX编写，PTX就是CUDA指令集架构，PTX参考手册中描述了PTX。通常PTX效率高于像C一样的高级语言。无论是使用PTX还是高级语言，内核都必须使用nvcc编译成二进制代码才能在设备上执行。
对于具体的利用nvcc如何编译及参数可参看C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.0\doc\pdf 的CUDA_Compiler_Driver_NVCC.pdf文件。


Create Kernels from CU Files
Run the Kernel
Determine Input and Output Correspondence
Kernel Object Properties
Specify Entry Points
Provide C Prototype Input
Complete Kernel Workflow
在.cu文件中创建 kernels

介绍如何含有kernel的cu文件生成PTX文件(parallelthread execution)

Compile a PTX File
如果你已经有一个cu文件myfun.cu,你想在GPU上执行，第一你必须将其编译生成PTX文件，一种方式就是利用CUDA toolkit的nvcc编译器，你可以在命令行窗口输入如下命令生成PTX文件

nvcc -ptx myfun.cu产生文件为myfun.ptx.
在matlab command输入！nvcc -ptx myfun.cu
Construct the Kernel Object
利用.cu文件盒.ptx文件你就可以在matlab中创建一个kernel 对象，即可以使用kernel计算：

k = parallel.gpu.CUDAKernel(‘myfun.ptx’, ‘myfun.cu’);

• 
  
  
  

  Note  You cannot save or load kernelobjects.

Run the Kernel
使用feval函数去在GPU上执行kernel .下面将介绍如何用gpuArray对象和MATLAB workspace data去执行kernel

Use Workspace Data
假设你已经有个kernel去执行两个向量的卷积，你就可以使用rand生成输入向量执行：
k = parallel.gpu.CUDAKernel(‘conv.ptx’, ‘conv.cu’);o = feval(k, rand(100, 1), rand(100, 1));即使输入是matlab 的workspace 数据（cpu数据），输出的也将是gpuArray类型


Use GPU Data
如果直接使用gpuArray作为输入将会更有效：
k = parallel.gpu.CUDAKernel(‘conv.ptx’, ‘conv.cu’);i1 = gpuArray(rand(100, 1, ‘single’));i2 = gpuArray(rand(100, 1, ‘single’));o1 = feval(k, i1, i2);因为输出是 gpuArray, 你可以向其他gpuArray数据一样去执行GPU操作，最终使用gather返回CPU数据到workspace
o2 = feval(k, o1, i2);r1 = gather(o1);r2 = gather(o2);
Determine Input and Output Correspondence
当调用[out1, out2] = feval(kernel, in1, in2,in3)时, 输入 in1, in2,and in3会对应于在cu文件中的c函数的输入参数 ，输出 out1 and out2存储当kernel执行后的第一第二C函数输入的非常量指针参数值。

例 cu文件中声明如下kernel:
void reallySimple( float * pInOut, float c )对应的matlab kernel object (k) 有如下属性:
MaxNumLHSArguments: 1   
NumRHSArguments: 2     
ArgumentTypes: {‘inout single vector’  ‘in single scalar’}
因此要执行feval,必须提供两个输入参数，可以有一个输出
y = feval(k, x1, x2)输入 x1 ，x2 对应C 函数原型的 pInOut和 c 参数. 输出y对应于 C kernel执行后的 pInOut值。

下面的例子将会复杂一些结合了常量和非常量指针：
void moreComplicated( const float * pIn, float * pInOut1, float * pInOut2 )对应matlab kernel属性:
MaxNumLHSArguments: 2   
NumRHSArguments: 3     
ArgumentTypes: {‘in single vector’  ‘inout single vector’  ‘inout single vector’}
可以使用 feval 在kernel (k)上执行：
[y1, y2] = feval(k, x1, x2, x3)三个输入参数x1,x2,x3对应于三个C 的输入参数，输出  y1， y2, 对应于kernel执行后的 pInOut1 and pInOut2 的值

Kernel Object Properties
当你创建一个内核对象没有加分号，或当你在命令行中键入对象变量， MATLAB显示内核对象的属性。
k = parallel.gpu.CUDAKernel(‘conv.ptx’, ‘conv.cu’)k =   parallel.gpu.CUDAKernel handle  Package: parallel.gpu  Properties:     ThreadBlockSize: [1 1 1]  MaxThreadsPerBlock: 512            GridSize: [1 1 1]    SharedMemorySize: 0          EntryPoint: ‘_Z8theEntryPf’  MaxNumLHSArguments: 1     NumRHSArguments: 2       ArgumentTypes: {‘in single vector’  ‘inout single vector’}一个内核对象的属性控制其执行行为。可使用.表示法来改变可以改变的那些属性。
可参看 CUDAKernel对象引用页。
Specify Entry Points
如果PTX文件包含多个入口点，您可以在myfun.ptx 中指定识别特定内核，你想要内核对象k指向：
k = parallel.gpu.CUDAKernel(‘myfun.ptx’, ‘myfun.cu’, ‘myKernel1’);
一个单一的PTX文件可以包含多个不同的内核的入口点。这些入口点中的每一个kernel都具有一个唯一的名称。
这些名称通常错位（如C + +中的重整） 。然而，当产生的NVCC PTX名称总是从CU包含原来的函数名。
如cu文件kernel定义为
__global__ void simplestKernelEver( float * x, float val )那么PTX 代码将会包含一个入口可能叫做 _Z18simplestKernelEverPff.
当你有多个入口点，在调用CUDAKernel时产生的内核时要为特定的内核指定项目名称

• 
  
  
  

  Note  The CUDAKernel function searches for yourentry name in the PTX file, and matches on any substring occurrences.Therefore, you should not name any of your entries as substrings of any others.

Provide C Prototype Input
如果没有PTX文件对应的CU 文件，你可以指定你的C内核来代替CU文件
k = parallel.gpu.CUDAKernel(‘myfun.ptx’, ‘float *, const float *, float’);C的原型语法支持C数据类型列在下面的表中。
[td]

Float Types	IntegerTypes	Boolean and Character Types
double, double2 float, float2	short, unsigned short, short2, ushort2 int, unsignedint, int2, uint2 long, unsigned long,  long2, ulong2 longlong,  unsigned long long,  longlong2, ulonglong2	bool char, unsignedchar, char2, uchar2

所有输入可以为标量或者指针也可以为 const.
一个内核的C声明的形式始终是：
__global__ void aKernel(inputs …)

• kernel 必须用void返回,操作只在输入参数 (scalars or pointers).
• kernel 不能分配任何内存，因此所有输出必须要在执行前预先分配好内存，所有输出大小必须在执行前就已知道
• 原则上，所有的指针传递到内核不能为const ，也可能包含输出数据，因为许多内核线程可能会修改该数据。
  
  

当把定义的kernel转成MATLAB时：

• 所有C输入标量 (double, float, int,etc.) 必须为MATLAB 标量, 或者 (i.e., single-element)gpuArray数据标量. 他们直接传递给kernel
• 所有 C输入常量指针 (constdouble *, etc.) 可以为MATLAB标量或者矩阵. 他们被转换为正确的类型拷贝到GPU上，指针的第一个元素被传递给kernel，原始大小不传递.  kernel直接接受结果通过 mxGetData在 mxArray上.
• 所有的C非常量指针输入以完全一样的非常指针转移到内核。然而，因为一个非常量指针可以改变，这将被作为一个内核的输出。
  

这些规则有一定的影响。最值得注意的是，一个内核每一个输出必然也是内核的输入，由于输入允许用户定义的输出的大小（如下将无法在GPU上分配内存） 。

Complete Kernel Workflow Add Two Numbers
这个例子在GPU进行两个双精度数据相加。

• cuda代码 .
  __global__ void add1( double * pi, double c ) {    *pi += c;}引导 __global__ 表面这是一个kernel入口，代码使用指针*pi返回输出结果，这是一个输入也是一个输出。将这个代码命名为test.cu在当前目录。
• 编译cu文件生成PTX 文件生成 test.ptx.
  nvcc -ptx test.cu
• 在matlab中创建如果有多个入口需要指定add1名字
  k = parallel.gpu.CUDAKernel(‘test.ptx’, ‘test.cu’);
• 执行两个输入1,1 默认一个线程
  >> o = feval(k, 1, 1);o =     2
  

Add Two Vectors
这个例子扩展了前面的两个矢量相加。为简单起见，假设向量中的元素等于一个线程块的线程的数量

• cu code
  
  1. __global__ void add2( double * v1, const double * v2 ) {    int idx = threadIdx.x;    v1[idx] += v2[idx];}
     

  复制代码
  //test.cu.
• Compile as before using nvcc.
  nvcc -ptx test.cu
• 指定入口名字
  k = parallel.gpu.CUDAKernel(‘test.ptx’, ‘add2’, ‘test.cu’);
• 设置正确的线程数
  >> o = feval(k, 1, 1);o =     2>> N = 128;>> k.ThreadBlockSize = N;>> o = feval(k, ones(N, 1), ones(N, 1));
  

参考文献：
http://www.mathworks.cn/cn/help/distcomp/executing-cuda-or-ptx-code-on-the-gpu.html
http://www.mathworks.cn/videos/introduction-to-gpu-computing-with-matlab-68770.html
http://www.mathworks.cn/cn/help/signal/examples/accelerating-correlation-with-gpus.html
http://www.mathworks.cn/cn/help/distcomp/using-gpuarray.html
http://www.mathworks.cn/discovery/matlab-gpu.html
http://www.mathworks.cn/company/newsletters/articles/gpu-programming-in-matlab.html

76361_92100v00_prototyping-algorithms-and-testing-cuda-kernels-in-matlab.pdf (2.06 MB, 下载次数: 531)