Roughly the same processing time for global and shared mem

[Ringworm]

Hello. I modified the simpleTexture program included in the SDK browser to perform a very simple blur. I’m using a block size of 15x15 and I’m returning an image without the borders (receive a 512x512 picture, return a 510x510).

NOT USING SHARED MEMORY: 738 ms

// calculate normalized texture coordinates

unsigned int x = blockIdx.x * blockDim.x + threadIdx.x + 1;

unsigned int y = blockIdx.y * blockDim.y + threadIdx.y + 1;

float u = x / (float) width;

float v = y / (float) height;

float stepu = 1 / (float) width;

float stepv = 1 / (float) height;

g_odata[(y - 1) * (width - 2) + (x - 1)] = (tex2D(tex, u + stepu, v) + tex2D(tex, u - stepu, v) + tex2D(tex, u, v + stepv) + tex2D(tex, u, v - stepv) + tex2D(tex, u + stepu, v + stepv) + tex2D(tex, u + stepu, v - stepv) + tex2D(tex, u - stepu, v + stepv) + tex2D(tex, u - stepu, v - stepv))/9;

USING SHARED MEMORY: 733 ms

shared float shared_data[BLOCK_SIZE][BLOCK_SIZE];

// calculate normalized texture coordinates

unsigned int x = blockIdx.x * blockDim.x + threadIdx.x + 1;

unsigned int y = blockIdx.y * blockDim.y + threadIdx.y + 1;

float u = x / (float) width;

float v = y / (float) height;

float stepu = 1 / (float) width;

float stepv = 1 / (float) height;



shared_data(ty, tx) = (tex2D(tex, u, v) + tex2D(tex, u + stepu, v) + tex2D(tex, u - stepu, v) + tex2D(tex, u, v + stepv) + tex2D(tex, u, v - stepv) + tex2D(tex, u + stepu, v + stepv) + tex2D(tex, u + stepu, v - stepv) + tex2D(tex, u - stepu, v + stepv) + tex2D(tex, u - stepu, v - stepv))/9;

__syncthreads();

g_odata[(y - 1) * (width - 2) + (x - 1)] = shared_data(ty, tx);

What’s wrong? Shouldn’t the texture mem → shared mem → global mem data transfers be faster than texture mem → global mem? How can I circumvent this problem?

Thank you.

[seibert]

What’s wrong? Shouldn’t the texture mem → shared mem → global mem data transfers be faster than texture mem → global mem?

No, there is no reason for that to be true. Shared memory is only an advantage if it allows you avoid unnecessary repetition of global memory reads, or if it helps you convert uncoalesced global memory reads into coalesced global memory reads. In the example you show, the texture cache is already doing this for you, so the storage in shared memory is completely unnecessary.

[wanderine]

What’s the size of the filter and what kind of GPU do you use? 730 ms sounds very slow for such a small image.

[Ringworm]

Uh, I think I might have forgotten to delete a loop calling the kernel. Now the processing time is about 611.9 ms.

My GPU has a compute capability of 1.1. Aside from the kernel, I’ve only modified this, from the simpleTexture example program:

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE, 1);

dim3 dimGrid((width - 2) / dimBlock.x, (height - 2) / dimBlock.y, 1);

What bemused me is that the processing time for the equivalent blur program executed only on the CPU was 11.7 ms. These were the matlab commands (the fourth line compiles the code):

I = imread(‘lena_bw.pgm’, ‘PGM’);

A = cast(I,‘double’);

A = (A/255);

nvmex -f nvmexopts.bat blurCPU.cu -IC:\cuda\include -LC:\cuda\lib -lcufft -lcudart

for n = 1:10; tic, blopOut = blurCPU(A);, tElapsedCPU(n) = toc; end; ExecOutCPU = mean(tElapsedCPU); clear tElapsedCPU; clear n;

Although I used the matcuda plugin for matlab on the last one, the difference is still appalling. Both programs’ codes are attached.

Previously, I had written two programs to denoise a 946x946 picture that called a function 100 times, both using matcuda and one of them 22x22 thread blocks. The one that was executed exclusively on the CPU ran 1.88 times more faster than the other (1311 ms vs 2461 ms). The memory allocation/deallocation and data transfer times were relatively close to each other (23 ms vs 32 ms), so it was really the denoise functions that made the difference. These were the matlab commands:

a=zeros(948,948); a(237:711,237:711)=1; b=a+0.2*randn(948,948); clear a;

for n = 1:10; tic, blopOut = outspock(b, 22, 0.011, 100);, tElapsedGPU(n) = toc; end; ExecOutGPU = mean(tElapsedGPU); clear tElapsedGPU; clear n; clear blopOut;

for n = 1:10; tic, blopCPU = outspockCPU(b, 0.011, 100);, tElapsedCPU(n) = toc; end; ExecCPU = mean(tElapsedCPU); clear tElapsedCPU; clear n; clear blopCPU;

Both are also attached (as outspockCPU.cu and outspock.cu).

Any help in what’s causing this will be appreciated.
blur_plus_denoise.rar (552 KB)

[wanderine]

Uh, I think I might have forgotten to delete a loop calling the kernel. Now the processing time is about 611.9 ms.

My GPU has a compute capability of 1.1. Aside from the kernel, I’ve only modified this, from the simpleTexture example program:

What bemused me is that the processing time for the equivalent blur program executed only on the CPU was 11.7 ms. These were the matlab commands (the third line compiles the code):

Although I used the matcuda plugin for matlab on the last one, the difference is still appalling. Both programs’ codes are attached.

Previously, I had written two programs to denoise a 946x946 picture that called a function 100 times, both using matcuda and one of them 22x22 thread blocks. The one that was executed exclusively on the CPU ran 1.88 times more faster than the other (1311 ms vs 2461 ms). The memory allocation/deallocation and data transfer times were relatively close to each other (23 ms vs 32 ms), so it was really the denoise functions that made the difference. These were the matlab commands:

Both are also attached (as outspockCPU.cu and outspock.cu).

Any help in what’s causing this will be appreciated.

You have to check the model of your GPU, if it does not have that many processor cores it can’t beat the CPU…

By the way, what is matcuda?

[Ringworm]

You have to check the model of your GPU, if it does not have that many processor cores it can’t beat the CPU…

By the way, what is matcuda?

It thought being a CUDA-enabled GPU always guaranteed some sort of speedup.

My CPU is an Intel Pentium D @ 3.0 GHz and I have 2GB of RAM.

These are my DeviceQuery results:

CUDA Device Query (Runtime API) version (CUDART static linking)

There is 1 device supporting CUDA

Device 0: “GeForce 8600 GT”

CUDA Driver Version: 3.0

CUDA Runtime Version: 3.0

CUDA Capability Major revision number: 1

CUDA Capability Minor revision number: 1

Total amount of global memory: 536543232 bytes

Number of multiprocessors: 4

Number of cores: 32

Total amount of constant memory: 65536 bytes

Total amount of shared memory per block: 16384 bytes

Total number of registers available per block: 8192

Warp size: 32

Maximum number of threads per block: 512

Maximum sizes of each dimension of a block: 512 x 512 x 64

Maximum sizes of each dimension of a grid: 65535 x 65535 x 1

Maximum memory pitch: 2147483647 bytes

Texture alignment: 256 bytes

Clock rate: 1.19 GHz

Concurrent copy and execution: Yes

Run time limit on kernels: Yes

Integrated: No

Support host page-locked memory mapping: No

Compute mode: Default (multiple host threads

can use this device simultaneously)

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 3.0, CUDA Runtime Versi

on = 3.0, NumDevs = 1, Device = GeForce 8600 GT

PASSED

Sorry, it’s not matcuda. It’s the MATLAB plug-in for CUDA, which you can find here: http://developer.nvidia.com/object/matlab_cuda.html. It uses the mex interface as a gateway to pass values through MATLAB to your C code. Matcuda , now called GPUmat, is a toolbox that allows standard MATLAB code to run on GPUs (http://gp-you.org/), but it’s still in an early stage of development.

[seibert]

It thought being a CUDA-enabled GPU always guaranteed some sort of speedup.

This is definitely not true. Although the 8600 GT is not the slowest CUDA-enabled GPU, it is pretty close. In addition, if your task is dominated by the time required to copy the data to the GPU, then no matter how fast your GPU is, the CPU could be faster.

It does sound like something odd is going on, but even once that is solved, it is quite possible that CUDA will still be slower than the CPU for this task.

[Ringworm]

This is definitely not true. Although the 8600 GT is not the slowest CUDA-enabled GPU, it is pretty close. In addition, if your task is dominated by the time required to copy the data to the GPU, then no matter how fast your GPU is, the CPU could be faster.

It does sound like something odd is going on, but even once that is solved, it is quite possible that CUDA will still be slower than the CPU for this task.

I see. But, as said, the data transfer is not what’s causing the bottleneck: «The memory allocation/deallocation and data transfer times were relatively close to each other (23 ms vs 32 ms), so it was really the denoise functions that made the difference». I guess I’ll have to try the tile division approach described here: http://www.nvidia.com/content/nvision2008/…o_with_CUDA.pdf.

[seibert]

And you are sure that this isn’t CUDA initialization included in the timing? Initializing the CUDA context can be kind of slow, and the first time you run a particular kernel, the driver does some further one-time initialization. Also, which OS is this? I believe the CUDA context initialization takes longer on Windows Vista/7, but I don’t have the posts handy.

[Ringworm]

And you are sure that this isn’t CUDA initialization included in the timing? Initializing the CUDA context can be kind of slow, and the first time you run a particular kernel, the driver does some further one-time initialization. Also, which OS is this? I believe the CUDA context initialization takes longer on Windows Vista/7, but I don’t have the posts handy.

Well, I didn’t use a warm-up kernel on the denoise programs, but I did on the blur. Or maybe I misunderstood your question.

I’m using Windows XP 32bits.

Also, at the “simple filter example” on the pdf from my last post, is g_idata on texture memory?