Ohhh, Ian, thanks :-) I confess to being very entry level with only a few days here and there at present. I've got a background in SMP, distributed computation and general multi-core work (and 8 bit machine code from way-back-when) but I've yet to read up on the CUDA architecture properly.
When I get a moment I'll plumb your code in and give it a go - 100x is more like what I was hoping for :-) i. On 29 January 2010 02:55, Ian Cullinan <[email protected]> wrote: > 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 > -- Ian Ozsvald (Professional Screencaster) [email protected] http://ProCasts.co.uk/examples.html http://TheScreencastingHandbook.com http://IanOzsvald.com + http://ShowMeDo.com http://twitter.com/ianozsvald
_______________________________________________ PyCUDA mailing list [email protected] http://host304.hostmonster.com/mailman/listinfo/pycuda_tiker.net
