PyCUDA March 2012

pycuda@tiker.net

18 participants
43 discussions

Re: [PyCUDA] curandom, access to generator

by Andreas Kloeckner

On Sat, 28 Jan 2012 18:21:29 -0500, Thomas Wiecki <Thomas_Wiecki(a)brown.edu> wrote: > I am currently revisiting this but having some problems with the > random number generator. > > generator.generators_per_block is 512 on my card, so I initialize 512 > generators, but I see that some of them don't produce random numbers > when sampling from them. I notice that in some subtle ways (mainly > that the distribution is not correct or all numbers are the same) if I > sample from more than 300-350 generators (always the last ones are > affected), but it's fine when using e.g. 128. So it seems I can only > use a smaller number of generators than what the card says I should be > able to use. > > Any idea on why that might be or how to investigate this further? Mysterious. What generator is this using? XORWOW? Tomasz, any ideas? Andreas

9 years, 2 months

Re: [PyCUDA] Now what?! I have pycuda installed and working but...

by Serra, Mr. Efren, Contractor, Code 7542

Excellent! That's the sort of response I was hoping for. Efren A. Serra (Contractor) DeVine Consulting, Inc. Naval Research Laboratory Marine Meteorology Division 7 Grace Hopper Ave., STOP 2 Monterey, CA 93943 Code 7542 Office: 831-656-4650 -----Original Message----- From: marmaduke [mailto:duke@eml.cc] Sent: Wednesday, March 28, 2012 9:37 AM To: Serra, Mr. Efren, Contractor, Code 7542 Cc: pycuda(a)tiker.net Subject: Re: [PyCUDA] Now what?! I have pycuda installed and working but... On Tue, Mar 27, 2012 at 22:34, Serra, Mr. Efren, Contractor, Code 7542 <efren.serra.ctr(a)nrlmry.navy.mil> wrote: > That's the sort of coding samples I haven't found online. PyCUDA (in my code at least) takes all/most of the pain out of GPU computing, but it's not meant to replace CUDA. Thus, if you search for CUDA examples, you will get much further with respect to understanding why/how to use CUDA. Then, you should work with those examples from PyCUDA. I would strongly suggest starting with NVIDIA's CUDA Programming Guide which is provided with the SDK.

9 years, 4 months

Now what?! I have pycuda installed and working but...

by Serra, Mr. Efren, Contractor, Code 7542

Trying to figure out where to boost performance for my apps. For instance, where should one use PyCUDA, attached is a program I'm trying to optimize or improve runtime. I, naively, started optimizing numpy.stdev and numpy.average. Is that what I should focus on? Thanks for your time. Efren A. Serra (Contractor) DeVine Consulting, Inc. Naval Research Laboratory Marine Meteorology Division 7 Grace Hopper Ave., STOP 2 Monterey, CA 93943 Code 7542 Office: 831-656-4650

9 years, 4 months

Re: [PyCUDA] POW support for Complex Types

by Andreas Kloeckner

<#part sign=pgpmime> On Tue, 27 Mar 2012 14:30:44 -0400, Alex Nitz <alex.nitz(a)ligo.org> wrote: > It seems that pycuda does not currently support the ** operator (pow) for > complex dtypes (complex64, complex128). Is this something that might be on > the horizon? > > I am currently working on a project that would benefit from the pycuda > gpuarray working the same as the numpy array. We would like to accelerate > portions of our code, but need to maintain full compatibility on machines > that do not have a GPU. How alike is the gpuarray intended to become to > the numpy array? > > I understand that this may not be on anyone else's priority list, so I may > try to make a stab at a patch for this if nobody else comes along. I don't > have much experience with the inner workings of pycuda yet, so that may > take a while, but if I get something together, can someone enlighten me on > the patch process, or point me to a link? Huh, weird. pow(complex, complex) is implemented in the C++ header, so in principle this should work. Can you insert a keep=True in the call to get_elwise_kernel in get_pow_kernel in pycuda/elementwise.py? That should give you the code for the failing kernel, and a better shot at making sense of the nvcc error message. Andreas

9 years, 4 months

POW support for Complex Types

by Alex Nitz

It seems that pycuda does not currently support the ** operator (pow) for complex dtypes (complex64, complex128). Is this something that might be on the horizon? I am currently working on a project that would benefit from the pycuda gpuarray working the same as the numpy array. We would like to accelerate portions of our code, but need to maintain full compatibility on machines that do not have a GPU. How alike is the gpuarray intended to become to the numpy array? I understand that this may not be on anyone else's priority list, so I may try to make a stab at a patch for this if nobody else comes along. I don't have much experience with the inner workings of pycuda yet, so that may take a while, but if I get something together, can someone enlighten me on the patch process, or point me to a link? Thanks, - Alex Nitz I try something like, import pycuda.autoinit import pycuda.gpuarray import numpy cpu = numpy.array([1+1j],dtype=numpy.complex64) gpu = pycuda.gpuarray.to_gpu(cpu) gpu = gpu **2 and get this error Traceback (most recent call last): File "pc.py", line 9, in <module> gpu = gpu **2 File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/gpuarray.py", line 586, in __pow__ func = elementwise.get_pow_kernel(self.dtype) File "<string>", line 2, in get_pow_kernel File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/tools.py", line 404, in context_dependent_memoize result = func(*args) File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/elementwise.py", line 545, in get_pow_kernel "pow_method") File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/elementwise.py", line 165, in get_elwise_kernel arguments, operation, name, keep, options, **kwargs) File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/elementwise.py", line 151, in get_elwise_kernel_and_types keep, options, **kwargs) File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/elementwise.py", line 75, in get_elwise_module options=options, keep=keep) File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/compiler.py", line 282, in __init__ arch, code, cache_dir, include_dirs) File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/compiler.py", line 272, in compile return compile_plain(source, options, keep, nvcc, cache_dir) File "/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/compiler.py", line 129, in compile_plain cmdline, stdout=stdout, stderr=stderr) pycuda.driver.CompileError: nvcc compilation of /tmp/tmpePwmmN/kernel.cufailed [command: nvcc --cubin -arch sm_20 -I/home/ahnitz/local/pycuda/lib/python2.7/site-packages/pycuda-2011.2.2-py2.7-linux-x86_64.egg/pycuda/../include/pycuda kernel.cu] [stderr: kernel.cu(19): error: no suitable conversion function from "pycuda::complex<float>" to "float" exists kernel.cu(19): error: no suitable conversion function from "pycuda::complex<float>" to "float" exists 2 errors detected in the compilation of "/tmp/tmpxft_000036ca_00000000-4_kernel.cpp4.ii".

9 years, 4 months

Re: [PyCUDA] pointers on adding 3d grid support

by Andreas Kloeckner

<#part sign=pgpmime> On Sun, 25 Mar 2012 22:01:52 -0300, Ezequiel Alfíe <ealfie(a)gmail.com> wrote: > Hello, > > I know that pycuda currently supports 1d and 2d grids only. > > I read a previous post on that topic > http://www.mail-archive.com/pycuda@tiker.net/msg02448.html > > What is the "new style launch interface"? cuLaunchKernel http://goo.gl/fHIjM "old style" is cuLaunchGrid, cuLaunch > Any pointers on what changes should I do in order to support 3d grids? - First, expose cuLaunchKernel to Python. I'd prefer that as _launch_kernel (with the leading underscore, this will be a private interface in Python). #ifdef this to only be enabled on CUDA 4 and higher. - Then, hack pycuda/driver.py to make Function.__call__ and the .prepared*() functions target _launch_kernel when run on CUDA 4.0 or higher. Do this by setting the Function methods depending on the CUDA version. That way, the check for CUDA 4.0 happens once and does not incur extra invocation overhead. That should do the trick and keep us backward compatible as much as possible. Good luck, and let me know if you need further help. Andreas

9 years, 4 months

pointers on adding 3d grid support

by Ezequiel Alfíe

Hello, I know that pycuda currently supports 1d and 2d grids only. I read a previous post on that topic http://www.mail-archive.com/pycuda@tiker.net/msg02448.html What is the "new style launch interface"? Any pointers on what changes should I do in order to support 3d grids? Thanks, Ezequiel

9 years, 4 months

Re: [PyCUDA] Reserving gpu global memory and assigning to global symbol

by Ezequiel Alfíe

On Sat, Mar 24, 2012 at 20:02, Ezequiel Alfíe <ealfie(a)gmail.com> wrote: > Answering my own question: > > The code below seems to work but is odd enough to doubt about. > > I guess there's a better, like the intended, way to do it... do you > have any suggestions? > > ## file test.cu > > #include <stdio.h> > > extern "C" { > __global__ void use_my_array(); > __device__ float *my_array; > } > /* simple kernel just to see values of my_array */ > __global__ void use_my_array() { > int i = threadIdx.x; > printf("%.2e\n", my_array[i]); > return; > } > > ## compile with: nvcc -arch=sm_21 -cubin test.cu > > ## file test.py > > import struct > import numpy as np > import pycuda.autoinit > from pycuda.gpuarray import GPUArray > import pycuda.driver as drv > > mod=drv.module_from_file('test.cubin') > use_my_array=mod.get_function('use_my_array') > > #array of 4 float32s on host with random values > > array_host=np.zeros((4,), np.float32) > r=np.random.random(4) > array_host[:] = r[:] > > #allocate from pycuda > array_gpu=GPUArray((4,), np.float32) > > #copy values from array_host > array_gpu.set(array_host) > > # get pointer of pycuda allocated gpu array > gpu_pointer = array_gpu.ptr > > # get address of gpu_pointer and its size > my_array_pointer, pointer_size = mod.get_global('my_array') > > # convert gpu_pointer into string of bytes and > # copy it to the address of my_array pointer > drv.memcpy_htod(my_array_pointer, struct.pack("Q", long(gpu_pointer))) According to python's struct module, I guess this should be "@p" (native pointer) instead of "Q" (long int). (of course happen to be the same on amd64) > > use_my_array(block=(4,1,1)) > print array_host > > > On Sat, Mar 24, 2012 at 19:08, Ezequiel Alfíe <ealfie(a)gmail.com> wrote: >> Hello everyone. I'd want to know: >> >> is it possible (how?) to reserve memory from pycuda and assign the >> resulting device pointer to a "global" pointer on a cuda module? >> >> like, for instance, having inside a test.cu file >> >> __device__ float *my_array; >> >> __global__ void somekernel() { >> int i = threadIdx.x; >> my_array[i] = 0.3f; >> } >> >> (have test.cu compiled into cubin test.cubin) >> >> and from python calling: >> import pycuda.autoinit >> import pycuda.driver as drv >> mod=drv.module_from_file('test.cubin') >> >> #dynamcally choose size >> size = 4 * 1000 # 4 is sizeof float32 >> mem1 = drv.mem_alloc(size) >> mem1_pointer = (long) mem1 >> >> my_array_pointer = mod.get_global('my_array') >> >> memcpy_htod(my_array_pointer, mem1_pointer) >> # this fails with TypeError: expected a readable buffer object >> >> -- >> How do I write the mem1_pointer to my_array ? >> Is there a simpler aproach to reserving global memory from pycuda and >> assigning to a global pointer variable so it is usable from kernels? >> (without passing it as kernel parameter) >> >> Thanks, >> Ezequiel

9 years, 4 months

Re: [PyCUDA] Reserving gpu global memory and assigning to global symbol

by Ezequiel Alfíe

Answering my own question: The code below seems to work but is odd enough to doubt about. I guess there's a better, like the intended, way to do it... do you have any suggestions? ## file test.cu #include <stdio.h> extern "C" { __global__ void use_my_array(); __device__ float *my_array; } /* simple kernel just to see values of my_array */ __global__ void use_my_array() { int i = threadIdx.x; printf("%.2e\n", my_array[i]); return; } ## compile with: nvcc -arch=sm_21 -cubin test.cu ## file test.py import struct import numpy as np import pycuda.autoinit from pycuda.gpuarray import GPUArray import pycuda.driver as drv mod=drv.module_from_file('test.cubin') use_my_array=mod.get_function('use_my_array') #array of 4 float32s on host with random values array_host=np.zeros((4,), np.float32) r=np.random.random(4) array_host[:] = r[:] #allocate from pycuda array_gpu=GPUArray((4,), np.float32) #copy values from array_host array_gpu.set(array_host) # get pointer of pycuda allocated gpu array gpu_pointer = array_gpu.ptr # get address of gpu_pointer and its size my_array_pointer, pointer_size = mod.get_global('my_array') # convert gpu_pointer into string of bytes and # copy it to the address of my_array pointer drv.memcpy_htod(my_array_pointer, struct.pack("Q", long(gpu_pointer))) use_my_array(block=(4,1,1)) print array_host On Sat, Mar 24, 2012 at 19:08, Ezequiel Alfíe <ealfie(a)gmail.com> wrote: > Hello everyone. I'd want to know: > > is it possible (how?) to reserve memory from pycuda and assign the > resulting device pointer to a "global" pointer on a cuda module? > > like, for instance, having inside a test.cu file > > __device__ float *my_array; > > __global__ void somekernel() { > int i = threadIdx.x; > my_array[i] = 0.3f; > } > > (have test.cu compiled into cubin test.cubin) > > and from python calling: > import pycuda.autoinit > import pycuda.driver as drv > mod=drv.module_from_file('test.cubin') > > #dynamcally choose size > size = 4 * 1000 # 4 is sizeof float32 > mem1 = drv.mem_alloc(size) > mem1_pointer = (long) mem1 > > my_array_pointer = mod.get_global('my_array') > > memcpy_htod(my_array_pointer, mem1_pointer) > # this fails with TypeError: expected a readable buffer object > > -- > How do I write the mem1_pointer to my_array ? > Is there a simpler aproach to reserving global memory from pycuda and > assigning to a global pointer variable so it is usable from kernels? > (without passing it as kernel parameter) > > Thanks, > Ezequiel

9 years, 4 months

Reserving gpu global memory and assigning to global symbol

by Ezequiel Alfíe

Hello everyone. I'd want to know: is it possible (how?) to reserve memory from pycuda and assign the resulting device pointer to a "global" pointer on a cuda module? like, for instance, having inside a test.cu file __device__ float *my_array; __global__ void somekernel() { int i = threadIdx.x; my_array[i] = 0.3f; } (have test.cu compiled into cubin test.cubin) and from python calling: import pycuda.autoinit import pycuda.driver as drv mod=drv.module_from_file('test.cubin') #dynamcally choose size size = 4 * 1000 # 4 is sizeof float32 mem1 = drv.mem_alloc(size) mem1_pointer = (long) mem1 my_array_pointer = mod.get_global('my_array') memcpy_htod(my_array_pointer, mem1_pointer) # this fails with TypeError: expected a readable buffer object -- How do I write the mem1_pointer to my_array ? Is there a simpler aproach to reserving global memory from pycuda and assigning to a global pointer variable so it is usable from kernels? (without passing it as kernel parameter) Thanks, Ezequiel

9 years, 4 months

2021

2020

2019

2018

2017

2016

2015

2014

2013

2012

2011

2010

2009

2008

PyCUDA March 2012