You think that's a speedup? :P
You're only using of the multiprocessors in your GPU! (Because you're launching
a 1x1 grid). Try this on for size:
======
import pycuda.driver as drv
import pycuda.tools
import pycuda.autoinit
import numpy
import numpy.linalg as la
from pycuda.compiler import SourceModule
blocks = 64
block_size = 128
nbr_values = blocks * block_size
n_iter = 100000
#############
# GPU SECTION
mod = SourceModule("""
__global__ void addone(float *dest, float *a, int n_iter)
{
const int i = blockDim.x*blockIdx.x + threadIdx.x;
for(int n = 0; n < n_iter; n++) {
a[i] = sin(a[i]);
}
dest[i] = a[i];
}
""")
addone = mod.get_function("addone")
a = numpy.ones(nbr_values).astype(numpy.float32)
a += 1 # a is now an array of 2s
dest = numpy.zeros_like(a)
start = drv.Event()
end = drv.Event()
start.record()
addone(drv.Out(dest), drv.In(a), numpy.int32(n_iter), grid=(blocks,1),
block=(block_size,1,1))
#stop timer
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print "GPU time:", secs
print "GPU result starts with...", dest[:3]
#############
# CPU SECTION
a = numpy.ones(nbr_values).astype(numpy.float32)
a += 1 # a is now an array of 2s
start.record()
for i in range(n_iter):
a = numpy.sin(a)
#stop timer
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print "CPU time:", secs
print "CPU result starts with...", a[:3]
======
(I reduced the number of iterations so it doesn't take forever on the CPU). On
my machine (3GHz Core 2 Duo, GTX280, Linux), I get:
GPU time: 0.0843682250977
GPU result starts with... [ 0.00547702 0.00547702 0.00547702]
CPU time: 8.12050439453
CPU result starts with... [ 0.00547701 0.00547701 0.00547701]
So, about a 100x speedup for the GPU version. Would be more a speedup with more
iterations (the overhead of copying the data to the GPU and back is the same
regardless). In fact, for such a small amount of data (only 32K) you can
probably significantly increase the size of the data without incurring much
more copying overhead - setting up the transfer is expensive, copying a few KB
isn't.
Have a play with the params and enjoy.
Next thing to get even more speed is to copy the data to shared memory within
each block, do the computation there and then copy the result back to main
memory when you're done. The NVIDIA docs and whitepapers should make it fairly
clear how to achieve that :)
Cheers,
Ian Cullinan
________________________________________
From: [email protected] [[email protected]] On Behalf Of Ian
Ozsvald [[email protected]]
Sent: Friday, 29 January 2010 1:29 AM
To: [email protected]
Subject: [PyCUDA] Very simple speed testing code for another beginner...
Here is some very simple speed testing code - maybe it is useful to another
beginner. I've used it to convince myself that GPUs really do go faster than
CPUs (this is useful here in the office to show my physics colleagues).
The code was adapted from hello_gpu.py. It has two halves, first it does a
simple calculation many times on the GPU and then it does the same calculation
on the CPU. Both times it uses dev.Event to count how long the operations took.
Roughly speaking on my WinXP Intel Core2 Duo 2.66GHz CPU (1 CPU used) the
9800GT GPU comes out 20-55* faster than the CPU.
In the code below a value for sin is calculated 2,000,000 times in a 400
element array. A 20-30* speedup holds for tan, sin, addition, sqrt, exp. The
pow function shows a 55* speedup. If you want to do your own testing then
replace the two references to 'sin' with your chosen function. Remember that
2,000,000 is also written twice so change it in both places to alter the number
of iterations. Extra note - the final result for 'tan' diverges quickly, 'sin'
and others seem to be mostly stable.
By lowering the iterations from 2,000,000 to 200 (in both places) then both the
CPU and GPU complete their tasks in roughly the same time.
I did a variation where 'dest' is removed and 'float *a' is referenced by
drv.InOut(a) (so a is the input parameter and is also used for the output
result) - I didn't observe any obvious speed difference.
Side note - I'm also using the NVIDIA System Monitor, I've selected all the GPU
outputs along with CPU outputs so they hover as transparent displays at the top
of the screen. Whenever the GPU is invoked you see the GPU Usage, Cooler and
Temp change.
HTHs another newbie,
Ian.
_______________________________________________
PyCUDA mailing list
[email protected]
http://host304.hostmonster.com/mailman/listinfo/pycuda_tiker.net
