PyCUDA July 2011

pycuda@tiker.net

17 participants
36 discussions

2011.1 released

by Andreas Kloeckner

Hi all, I've just released PyCUDA version 2011.1. This is, once again, a rather big release. A detailed list of changes is below. In the same spirit as PyOpenCL (whose 2011.1 release happened yesterday), I'll try to move to smaller, more frequent releases in the future. Have fun, and let me know if there are any issues. Andreas Detailed list of changes in 2011.1: * Add support for CUDA 3.0-style OpenGL interop. (thanks to Tomasz Rybak) * Add pycuda.driver.Stream.wait_for_event(). * Add range and slice keyword argument to pycuda.elementwise.ElementwiseKernel.__call__(). * Document preamble constructor keyword argument to pycuda.elementwise.ElementwiseKernel. * Add vector types, see pycuda.gpuarray.vec. * Add pycuda.scan. * Add support for new features in CUDA 4.0. * Add pycuda.gpuarray.GPUArray.strides, pycuda.gpuarray.GPUArray.flags. Allow the creation of arrys in C and Fortran order. * Adopt stateless launch interface from CUDA, deprecate old one. * Add CURAND wrapper. (with work by Tomasz Rybak) * Add pycuda.compiler.DEFAULT_NVCC_FLAGS.

8 years, 9 months

Custom reduction using parameters in reduce expressions

by Ryan Marcus

Hello! I'm a bit of a pycuda newb, and today I dove into the source for the first time. Essentially, I was trying to implement a numpy-style argmax argument using a custom reduction kernel: http://docs.scipy.org/doc/numpy/reference/generated/numpy.argmax.html I thought this would be a pretty simple procedure, and the code I tried was: maxloc_reduction_k = ReductionKernel( numpy.int32, arguments="float *x", neutral="0", map_expr="i", reduce_expr="(x[(int)a] > x[(int)b]) ? (int)a : (int)b") What I discovered was that pycuda really didn't appreciate me using the specified arguments (in this case, "float *x") in the reduce_expr part. Looking at reduction.py (in the latest release code, 0.94.2) , I couldn't really see a reason for why this shouldn't be possible. I added my own parameter to the reduction kernel's constructor ("hack") to allow my original parameters to be passed to both _stage1 and _stage2. This allowed me to do this: maxloc_reduction_k = ReductionKernel( numpy.int32, hack=True, arguments="float *x", neutral="0", map_expr="i", reduce_expr="(x[(int)a] > x[(int)b]) ? (int)a : (int)b") This way, my hacked changes would only affect some reduction kernels. I'm not totally clear on how to create a patch. Hopefully I didn't screw it up. I don't really know if I'm allowed to send attachments to a mailing group... so I used pastebin. Here's my changes: http://pastebin.com/R5i5JveM Just in case I messed up creating a patch, here's my full modified version of reduction.py: http://pastebin.com/WExgBTQ9 Is there a better way to be doing this? I considered using the current dev-build of pycuda to access float2, and store my array value in x and my index in y, but the reduction code hasn't been updated to properly handle float2's (you can't just assert a float to another volatile type -- you have to copy the .x and the .y). If this is functionality that is actually missing, I'll try to actually do the work to implement this in a non-hackish way. Thanks, Ryan Marcus

8 years, 10 months

Error in NVCC compilation when options are given

by Scott Biersdorff

Hi, I would like to report a possible bug with pycuda. To show the issue I've modified demo.py adding by passing a option to the nvcc compiler: mod = SourceModule(""" __global__ void doublify(float *a) { int idx = threadIdx.x + threadIdx.y*4; a[idx] *= 2; } """, options=['-ccbin /usr/bin']) However I get the following error when running demo.py: python demo.py *** compiler options is ['-ccbin /usr/bin', '-arch', 'sm_12', '-I/usr/local/lib/python2.7/dist-packages/pycuda-0.94.2-py2.7-linux-x86_64.egg/pycuda/../include/pycuda'] Traceback (most recent call last): File "demo.py", line 22, in <module> """, options=['-ccbin /usr/bin']) File "/usr/local/lib/python2.7/dist-packages/pycuda-0.94.2-py2.7-linux-x86_64.egg/pycuda/compiler.py", line 239, in __init__ arch, code, cache_dir, include_dirs) File "/usr/local/lib/python2.7/dist-packages/pycuda-0.94.2-py2.7-linux-x86_64.egg/pycuda/compiler.py", line 229, in compile return compile_plain(source, options, keep, nvcc, cache_dir) File "/usr/local/lib/python2.7/dist-packages/pycuda-0.94.2-py2.7-linux-x86_64.egg/pycuda/compiler.py", line 108, in compile_plain cmdline, stdout=stdout, stderr=stderr) pycuda.driver.CompileError: nvcc compilation of /tmp/tmp4keOYP/kernel.cu failed [command: nvcc --cubin -ccbin /usr/bin -arch sm_12 -I/usr/local/lib/python2.7/dist-packages/pycuda-0.94.2-py2.7-linux-x86_64.egg/pycuda/../include/pycuda kernel.cu] [stderr: nvcc fatal : Unknown option 'ccbin /usr/bin' ] Notice that the nvcc command listed in the output is valid and will compile the kernel correctly, yet I receiving an error from pycuda. I would welcome any suggestions or workarounds, I'm using pycuda 0.94.2 from source with python 2.7.1+ on Ubuntu 11.04. Thanks, - Scott

8 years, 11 months

New occupancy values - patch

by Tomasz Rybak

Hello. I attach patch updating pycuda.tools.DeviceData and pycuda.tools.OccupancyRecord to take new devices into consideration. I have tried to maintain "style" of those classes and introduced changes only when necessary. I have done changes using my old notes and NVIDIA Occupancy Calculator. Unfortunately I currently do not have access to Fermi to test those fully. Best regards. Tomasz Rybak

9 years, 1 month

How do I diagnose a CUDA launch failure due to being out of resources?

by Eli Stevens (Gmail)

This is a sorta-cross-post from stack overflow: http://stackoverflow.com/questions/6892280/how-do-i-diagnose-a-cuda-launch-… I'm getting an out-of-resources error when trying to launch a CUDA kernel (through PyCUDA), and I'm wondering if it's possible to get the system to tell me which resource it is that I'm short on. Obviously the system knows what resource has been exhausted, I just want to query that as well. I've used the occupancy calculator, and everything seems okay, so either there's a corner case not covered, or I'm using it wrong. I know it's not registers (which seems to be the usual culprit) because I'm using < 70 and it still fails with a 1x1x1 block and 1x1 grid on a CC 2.1 device. Thanks for any help. I posted a thread on the NVidia boards: http://forums.nvidia.com/index.php?showtopic=206261&st=0 But got no responses. If the answer is "you can't ask the system for that information" that would be nice to know too (sort of... ;). Thanks, Eli

9 years, 1 month

Re: [PyCUDA] trying to create a 3d texture

by Andreas Kloeckner

On Tue, 12 Jul 2011 19:06:52 -0700, Eli Stevens (Gmail) wrote: > Followup: > > What is class pycuda.driver.Memcpy3D.src_height supposed to be set > to? > I can't seem to find a valid value. See test_3d_texture() in test/test_wrapper.py. Andreas

9 years, 1 month

Changing nvcc command in configure.py at install time?

by Ben Aylott

Hi, Im using pycuda on arch linux with CUDA 4 and have had problems with the gcc 4.4 requirement (with arch using gcc-4.6 as default and providing gcc-4.4 as a separate package). I am aware of many possible solutions (which all work for me) but it would be quite nice to change the gcc that nvcc uses via 'nvcc -ccbin=/usr/bin/gcc-4.4' etc. You can obviously pass this to the compile function in the wrapper but it would be nice to be able to add some options at install time (via configure.py etc.) so a PKGBUILD can configure pycuda to 'work out the box'. Im not sure if this is already possible or if it would be feasible? regards, Ben Aylott

9 years, 1 month

pycuda installation error on centos 5.6

by Rekhesh Mohan

Hi people.. Greetings! :) Anyone got pycuda installed on centos 5.6? All my attempts with various combinations of default gcc, gcc44, default python, and python 2.6 ends up with the following error during make: g++ -pthread -shared build/temp.linux-x86_64-2.6/src/cpp/cuda.o build/temp.linux -x86_64-2.6/src/cpp/bitlog.o build/temp.linux- x86_64-2.6/src/wrapper/wrap_cudadr v.o build/temp.linux-x86_64-2.6/src/wrapper/mempool.o build/temp.linux- x86_64-2. 6/bpl-subset/bpl_subset/libs/python/src/object/enum.o build/temp.linux- x86_64-2. 6/bpl-subset/bpl_subset/libs/python/src/object/stl_iterator.o build/temp.linux-x 86_64-2.6/bpl-subset/bpl_subset/libs/python/src/object/life_support.o build/temp .linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/object/function.o build/ temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/object/iterator.o bu ild/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/object/function_ doc_signature.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/sr c/object/inheritance.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/py thon/src/object/class.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/p ython/src/object/pickle_support.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_sub set/libs/python/src/converter/arg_to_python_base.o build/temp.linux- x86_64-2.6/b pl-subset/bpl_subset/libs/python/src/converter/from_python.o build/temp.linux-x8 6_64-2.6/bpl- subset/bpl_subset/libs/python/src/converter/builtin_converters.o bu ild/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/converter/type_i d.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/converter/ registry.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/lon g.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/import.o b uild/temp.linux-x86_64-2.6/bpl-subset/bpl_subset/libs/python/src/str.o build/tem p.linux-x86_64-2.6/bpl-subset/bpl_subset/libs/python/src/exec.o build/temp.linux -x86_64-2.6/bpl-subset/bpl_subset/libs/python/src/object_operators.o build/temp. linux-x86_64-2.6/bpl-subset/bpl_subset/libs/python/src/wrapper.o build/temp.linu x-x86_64-2.6/bpl-subset/bpl_subset/libs/python/src/tuple.o build/temp.linux-x86_ 64-2.6/bpl-subset/bpl_subset/libs/python/src/slice.o build/temp.linux- x86_64-2.6 /bpl-subset/bpl_subset/libs/python/src/errors.o build/temp.linux- x86_64-2.6/bpl- subset/bpl_subset/libs/python/src/module.o build/temp.linux- x86_64-2.6/bpl-subse t/bpl_subset/libs/python/src/numeric.o build/temp.linux-x86_64-2.6/bpl- subset/bp l_subset/libs/python/src/list.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subse t/libs/python/src/object_protocol.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_s ubset/libs/python/src/dict.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subset/l ibs/smart_ptr/src/sp_collector.o build/temp.linux-x86_64-2.6/bpl- subset/bpl_subs et/libs/smart_ptr/src/sp_debug_hooks.o build/temp.linux-x86_64-2.6/bpl- subset/bp l_subset/libs/thread/src/pthread/once.o build/temp.linux-x86_64-2.6/bpl- subset/b pl_subset/libs/thread/src/pthread/thread.o build/temp.linux- x86_64-2.6/src/wrapp er/wrap_curand.o -L/usr/lib64 -L/usr/lib64 -lcuda -lcurand -lpython2.6 - o build/ lib.linux-x86_64-2.6/pycuda/_driver.so /usr/bin/ld: cannot find -lcurand collect2: ld returned 1 exit status By the way, I have set CXXFLAGS = ['-DBOOST_PYTHON_NO_PY_SIGNATURES'] in siteconf.py, and I have tried with pycuda-2011.1.2 (latest download available on website) as well as a clone from git. Any suggestions? Thanks in anticipation.. --R

9 years, 1 month

Re: [PyCUDA] share context between threads

by Andreas Kloeckner

On Mon, 25 Jul 2011 17:42:10 -0400, Anthony LaTorre <tlatorre9(a)gmail.com> wrote: Non-text part: multipart/mixed Non-text part: multipart/alternative > As of CUDA 4.0 it is possible to share a CUDA context across threads; is it > also possible to do this in the latest release of pyCUDA? If this is possible with the old-style push/pop API, then yes, otherwise no. Since I want PyCUDA to remain backward-compatible for a few releases of CUDA, I haven't switched it to CUDA 4's new context management. HTH, Andreas

9 years, 1 month

share context between threads

by Anthony LaTorre

As of CUDA 4.0 it is possible to share a CUDA context across threads; is it also possible to do this in the latest release of pyCUDA? Thanks, Tony

9 years, 1 month

2020

2019

2018

2017

2016

2015

2014

2013

2012

2011

2010

2009

2008

PyCUDA July 2011