Matlab GPU.docx
- 文档编号:4042767
- 上传时间:2022-11-27
- 格式:DOCX
- 页数:20
- 大小:29.44KB
Matlab GPU.docx
《Matlab GPU.docx》由会员分享,可在线阅读,更多相关《Matlab GPU.docx(20页珍藏版)》请在冰豆网上搜索。
MatlabGPU
MatlabGPU加速几步曲
GPU Powered Matlab
(一)
标签:
分类:
计算机与Internet
it matlabgpucuda并行运算
我的笔记本xialulee-pc3上面有一块ATI的显卡,但是由于A卡不支持CUDA,因此无法被Matlab的parralleltoolbox所用。
曾经考虑过用Aparapi写一段Java代码,然后在Matlab中通过Java接口来间接使用A卡的GPU,但在实施之前,我就把本本送给麻麻了。
因为没有电脑用,订购了一部便宜的台式机,上面装有一块GeForceGTX660。
虽不是什么高档卡,好在我对游戏的要求也不高,而且,当然,支持CUDA。
现在,我终于能够玩玩MatlabParallelToolbox提供的GPU计算功能了!
这里有两个问题。
其一,拿个什么样的程序作为Helloworld呢?
这个我以前就想好了,之前大战pythonchallege时,Level31印象尤为深刻(链接
function level31
% pythonchallenge level31
% 2012.03.18 AM 00:
05
% xialulee
% load image
[I, map] = imread('mandelbrot.gif');
[im_h, im_w] = size(I);
% set parameters
width = 0.036;
height = 0.027;
x_min = 0.34;
y_min = 0.57;
x_r = width / im_w;
y_r = height / im_h;
max_iter = 128;
tic
Z0 = bsxfun(@plus, ...
linspace(x_min, x_min+width-x_r, im_w), ...
1j * linspace(y_min+height-y_r, y_min, im_h).' ...
);
function output = calcidx(input)
% calc index value for each element in Z0
z = 0;
for k = 0 :
max_iter-1
z = z * z + input;
if abs(z) > 2, break; end
end
output = k;
end
idx = arrayfun(@calcidx, Z0);
toc
figure;
imshow(idx, map);
deltas = (idx - double(I)).';
deltas(deltas == 0) = [];
deltas(deltas == +16) = 0;
deltas(deltas == -16) = 1;
deltas = reshape(deltas, [23, 73]);
figure;
imshow(imresize(deltas.', 10, 'box'));
end
看看其中矩阵Z0的构造和idx的计算,我觉得都挺适合并行化的。
就拿Level31开刀,试试GPUPoweredMatlab到底有多厉害。
问题一解决。
问题二:
如何在Matlab中使用GPU?
之前没有做过,同时完全的CUDA零经验。
粗略看看MatlabParallelToolbox,意外地发现好像并不难。
立刻开工,开始对之前的level31.m进行修改,意外地发现只做了极少的修改,代码速度就获得了极大的提升!
代码如下:
function level31_gpu
% pythonchallenge level31
% GPU Powered
%
% 2013.03.24 PM 09:
08
% xialulee
% load image
[I, map] = imread('mandelbrot.gif');
[im_h, im_w] = size(I);
% set parameters
width = 0.036;
height = 0.027;
x_min = 0.34;
y_min = 0.57;
x_r = width / im_w;
y_r = height / im_h;
max_iter = 128;
tic
Z0 = bsxfun(@plus, ...
gpuArray.linspace(x_min, x_min+width-x_r, im_w), ...
1j * gpuArray.linspace(y_min+height-y_r, y_min, im_h).' ...
); % Create matrix Z0 and put it in the memory of the GPU
function output = calcidx(input)
% calc index value for each element in Z0
z = complex(0, 0); % For GPU. Declare that z is a complex number.
k = 0; % For GPU. Initialize k, or there will be a error
for k = 0 :
max_iter-1
z = z * z + input;
if abs(z) > 2, break; end
end
output = k;
end
idx = arrayfun(@calcidx, Z0);
idx = gather(idx); % Data transfer. GPU's => CPU's
toc
figure;
imshow(idx, map);
deltas = (idx - double(I)).';
deltas(deltas == 0) = [];
deltas(deltas == +16) = 0;
deltas(deltas == -16) = 1;
deltas = reshape(deltas, [23, 73]);
figure;
imshow(imresize(deltas.', 10, 'box'));
end
与之前的代码相比,有如下的几处改动
一、构造矩阵Z0时,使用gpuArray的静态方法linspace在显存上创建了数组。
不需要对bsxfun进行修改,因为一旦参数包含gpuArray类型,Matlab会自动使用对应的bsxfun。
二、在calcidx中,对k和z进行了初始化。
如果不初始化k,会报错:
ErrorusinggpuArray/arrayfun
Useofapotentiallyuninitializedvariable(s):
k erroratline:
37
如果不使用complex函数初始化z,会报错:
ErrorusinggpuArray/arrayfun
Variable'z'changedcomplexity.ConsiderdeclaringitCOMPLEXbeforelooporif-elsebody.erroratline:
35
由于这里的calcidx会被arrayfun施加到gpuArray的实例上,因此它再也不是一个普通的Matlab函数了,而是一个kernel。
我猜Matlab的ParralelToolbox采用了类似Aparapi的机制,由字节码反推AST,再遍历AST生成可被GPU执行的代码。
三、计算完成后,使用gather将显存中的结果导入到内存中
由于普通的,和GPU加速的代码都加了计时代码tic-toc,看看效果。
我现在很紧张,因为之前听说过很多GPU加速大失败的案例,而且我也没什么这方面的经验,虽然之前在《玩了玩ATIStream》小试过一把,不过失败了。
几次level31.m的运行结果
>>clearall
>>level31
Elapsedtimeis1.705742seconds.
>>level31
Elapsedtimeis1.711330seconds.
>>level31
Elapsedtimeis1.730576seconds.
>>level31
Elapsedtimeis1.707733seconds.
>>level31
Elapsedtimeis1.717991seconds.
再看看level31_gpu.m的效果
>>level31_gpu
Elapsedtimeis1.029675seconds.
>>level31_gpu
Elapsedtimeis0.021065seconds.
>>level31_gpu
Elapsedtimeis0.021019seconds.
>>level31_gpu
Elapsedtimeis0.021232seconds.
>>level31_gpu
Elapsedtimeis0.021126seconds.
第一次不怎么快(应该是编译消耗了不少时间),之后,和CPU版的level31.m相比,它就像是注射了类固醇的种马一样。
忘了说一句,CPU是i7-3770。
GPU的相关信息如下
>>gpuDevice
ans=
CUDADevicewithproperties:
Name:
'GeForceGTX660'
Index:
1
ComputeCapability:
'3.0'
SupportsDouble:
1
DriverVersion:
5
ToolkitVersion:
5
MaxThreadsPerBlock:
1024
MaxShmemPerBlock:
49152
MaxThreadBlockSize:
[1024102464]
MaxGridSize:
[2.1475e+096553565535]
SIMDWidth:
32
TotalMemory:
1.6106e+09
FreeMemory:
1.3412e+09
MultiprocessorCount:
6
ClockRateKHz:
888500
ComputeMode:
'Default'
GPUOverlapsTransfers:
1
KernelExecutionTimeout:
1
CanMapHostMemory:
1
DeviceSupported:
1
DeviceSelected:
1
没有对之前的level31.m做什么重大的修改,就实现了GPU加速,这真是一件令人心旷神怡的事情。
GPU Powered Matlab
(二)
标签:
itmatlabgpupythonchallenge
分类:
计算机与Internet
《GPUPoweredMatlab
(一)》中,第一次试验了在Matlab中使用GPU计算,结果非常令人满意,仅仅对代码进行了细微的修改,就实现了大约80倍的提速。
所以今天又小小的尝试了一番。
用什么作为试验的任务呢?
还是想起了PythonChallenge,因为其中包含了不少图像处理的关卡。
就拿Level16开刀吧。
这一关有一副图像,需要根据粉红色块的位置,对每一行的像素进行循环移动,才能得到有意义的图像。
我的设想是这样的:
用一个线程处理一行,由于图像有480行,所以需要480个线程。
线程内部实现粉红色块的搜索和像素的循环移动。
由于这一关的运算量并不大,因此我不期望使用GPU会有多少速度提升,只是想验证一下我设想的代码是否能像期望中那样运行。
由于每个线程的处理稍微有些复杂,似乎无法使用纯Matlab来描述,所以我试着使用CUDAC描述了对图像每一行的操作,如下
__global__ void rotpixel(const unsigned char * imin, unsigned char * imout, int w, int h) {
int row_num = threadIdx.x;
int first_pink = 0;
if (row_num < h) {
for (int n=0; n if (imin[w*row_num+n] == 195) { first_pink = n; break; } } if (first_pink) { for (int n=0; n imout[w*row_num+n] = imin[w*row_num+n+first_pink]; for (int n=0; n imout[w*row_num+w-first_pink+n] = imin[w*row_num+n]; } } } 在Matlab的CommandWindow中编译上面这段代码 >>! nvcc-ptxlevel16.cu 构造CUDAKernel对象 >>ko=parallel.gpu.CUDAKernel('level16.ptx','level16.cu'); 构造的时候为什么同时需要编译得到的ptx文件和源代码文件呢? 我猜大概是由于ptx没有提供足够的元信息,比如参数个数,参数类型等,Matlab需要通过分析源代码来得到这些必需的信息。 看看ko的属性 >>ko ko= CUDAKernelwithproperties: ThreadBlockSize: [111] MaxThreadsPerBlock: 1024 GridSize: [111] SharedMemorySize: 0 EntryPoint: '_Z8rotpixelPhS_ii' MaxNumLHSArguments: 1 NumRHSArguments: 4 ArgumentTypes: {'inuint8vector' 'inoutuint8vector' 'inint32scalar' 'inint32scalar'} 主要ArgumentTypes这个属性。 由于使用了const修饰第一个参数,所以它就是in参数,而第二个参数没有使用const,它是inout。 非指针类型的参数大抵都是in类型了。 下面是使用这个kernel的Matlab通关代码: function level16_gpu % level 16 % GPU Powered % 2013.03.25 PM 07: 06 % xialulee persistent ko; [I, cmap] = imread('mozart.gif'); if isempty(ko) ko = parallel.gpu.CUDAKernel('level16.ptx', 'level16.cu'); end tic [h, w] = size(I); ko.ThreadBlockSize = h; imout = gpuArray.zeros(w, h, 'uint8'); imout = feval_r(ko, I', imout, w, h); toc imshow(imout', cmap); end 值得注意的是代码中需要人工指定ThreadBlockSize,这一点在文档中有相应的说明。 代码中使用了tic-toc计时,看看加速效果 >>level16_gpu Elapsedtimeis0.001373seconds. >>level16_gpu Elapsedtimeis0.000951seconds. >>level16_gpu Elapsedtimeis0.000435seconds. >>level16_gpu Elapsedtimeis0.000436seconds. >>level16_gpu Elapsedtimeis0.000448seconds. 一如既往的第一次有点慢。 作为比较,将之前《xialulee大战pythonchallenge——level16》中的代码加入了tic-toc,如下 function level16 % level 16 % 2012.03.03 AM 10: 34 % xialulee [I, cmap] = imread('mozart.gif'); tic [h, w] = size(I); Ic = mat2cell(I, ones(1, h), w); I2 = cellfun(@rotpixels, Ic, 'UniformOutput', false); I3 = cell2mat(I2); toc imshow(I3, cmap); end function line = rotpixels( line ) idx = find(line==195, 1, 'first'); if idx line = [line(idx+1: end), line(1: idx)]; end end 看看时间消耗 >>level16 Elapsedtimeis0.021690seconds. >>level16 Elapsedtimeis0.013120seconds. >>level16 Elapsedtimeis0.012727seconds. >>level16 Elapsedtimeis0.012552seconds. >>level16 Elapsedtimeis0.013115seconds. 通过比较发现使用GPU确实比之前的快了。 符合期望,不过之前的代码并不能反映CPU的真实水准,构造cell和cellfun什么的做法估计效率是比较低下的。 反正这一次试验的目的不在于加速,只是想增加一些对GPU的了解。 GPU Powered Matlab(三)——关于Matlab的eps函数 标签: Matlabcudagpu相对浮点精度 分类: 计算机与Internet 为了了解eps函数的运作方式,我们先后用MATLAB的typecast,C/C++的位段,Pythonctypes的位段,以及Java提供的ulp实现了相对浮点精度的运算。 但是Matlab自身究竟是如何实现eps的我们依然无从得之,因为它不是开源的。 有时候想看看Octave是如何实现的,结果我翻了半天Octave的C++代码,什么也没找着。 今天突然想到,可以看看MATLABCoder会将eps函数翻译成怎样的C代码,于是编写了如下的Matlab函数: function y = testeps(x) % 2012.04.03 % xialulee assert(isa(x, 'double')); % y = eps(x); y = eps(x); end 因为Matlab和C分别属于动态类型和静态类型语言,为了解决冲突,MATLABCoder要求使用assert的方式来指定变量的类型,只有知道了类型才能生成相应的C代码。 使用codegen函数进行翻译,由于翻译的过程中会保留原Matlab代码中的注释,所以可以很容易找到“y=eps(x)”对应的C代码的位置。 看看,发现没有什么稀奇的,既没有用普通的位操作,也没有使用位段,而是使用了 frexp和ldexp。 虽然《关于Matlab的eps函数》系列中已经解释过算法的原理好多次,这里还是再简要提一句吧。 对于64位的754浮点数,其在内存中的布局为: SEEEEEEEEEEEMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMMM 符号位1位,带偏移的指数域11位,尾数位52位。 其代表的数值由下式计算: 1.M×2E-1023 其中1023是偏移量。 计算相对浮点精度很简单,将所有M置0,将指数减去52(52为尾数部分的长度),即可。 利用之前在《在Matlab中方便地操
- 配套讲稿:
如PPT文件的首页显示word图标,表示该PPT已包含配套word讲稿。双击word图标可打开word文档。
- 特殊限制:
部分文档作品中含有的国旗、国徽等图片,仅作为作品整体效果示例展示,禁止商用。设计者仅对作品中独创性部分享有著作权。
- 关 键 词:
- Matlab GPU