very simple cuda ptx code memory speed

3 ビュー (過去 30 日間)
Gaszton
Gaszton 2011 年 5 月 19 日
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

採用された回答

Edric Ellis
Edric Ellis 2011 年 5 月 20 日
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 件)

カテゴリ

Help Center および File ExchangeGPU Computing についてさらに検索

Community Treasure Hunt

Find the treasures in MATLAB Central and discover how the community can help you!

Start Hunting!

Translated by