MATLAB Answers

Gaszton
0

very simple cuda ptx code memory speed

Asked by Gaszton
on 19 May 2011
Hello, i made a simple cuda kernel to measure global memory transfer speed to the cuda processors:
__global__ void SR2add(float* dataout,const float* datain,int size) {
int mindex=blockIdx.x*blockDim.x+threadIdx.x;
if (mindex>=size)
return;
dataout[mindex]=datain[mindex];
}
The matlab function i wrote for it:
function GPU_MemBandTest()
import parallel.gpu.GPUArray
xsize=1024;
ysize=768;
vectorsize=xsize*ysize;
threadpblock=1024;
k=parallel.gpu.CUDAKernel('MemBandTest.ptx', 'MemBandTest.cu');
k.ThreadBlockSize=[threadpblock,1,1];
k.GridSize=[ceil(vectorsize/threadpblock),1];
ddatain=parallel.gpu.GPUArray.zeros(vectorsize,1,'single');
dataout=rand(vectorsize,1,'single');
ddataout=GPUArray(dataout);
tic
for i=1:1000
[ddataout]=feval(k,ddataout,ddatain,vectorsize);
end
time=toc;
disp(['ms time= ' num2str(time)])
disp([num2str(vectorsize*4/(time*10^6)) 'GB/s'])
end
I got ms time= 0.73629 and 4.2724GB/s result for that. I would like to ask: 1; that am i doing correctly the measurement? 2; Is there anything i can do to speed up this simple code or this is an expectable result for this kernel in matlab?
I have MATLAB R2011a, CUDA Toolkit 3.2, gt425m device, newest driver installed for it
If I use float* datain instead of const float* datain, the execution time goes up to 2.4ms
3; What could be the explanation of this?
Thanks for anyone who helps,
Gaszton

  0 Comments

Sign in to comment.

1 Answer

Answer by Edric Ellis
on 20 May 2011
 Accepted Answer

To answer your questions:
  1. You're a factor of 2 out in your bandwidth calculation because you need to consider that the data is being read and written by the device. So the total data transfer is twice the size of the data
  2. On my machine, with that factor of 2, I get 40 GB/s on a C2070, which is fairly reasonable. As the size of the data increases, this rate increases.
  3. If you make datain be non-const, we treat that as an input-output variable (you can see this from the properties of the CUDAKernel) and allocate space to store the result. So, for CUDAKernel, const-correctness is actually very important for performance!

  0 Comments

Sign in to comment.