PyCUDA March 2014

pycuda@tiker.net

5 participants
9 discussions

by Tomasz Rybak

Hello. I've been packaging PyCUDA for Debian. I run all the tests to ensure that package works on Python 2 and Python 3. All tests pass except for on from test_driver.py: $ python test_driver.py ============================= test session starts ============================== platform linux2 -- Python 2.7.5 -- pytest-2.3.5 collected 21 items test_driver.py ........F............ =================================== FAILURES =================================== _____________________ TestDriver.test_register_host_memory _____________________ args = (<test_driver.TestDriver instance at 0x24e7d88>,), kwargs = {} pycuda = <module 'pycuda' from '/usr/lib/python2.7/dist-packages/pycuda/__init__.pyc'> ctx = <pycuda._driver.Context object at 0x2504488> clear_context_caches = <function clear_context_caches at 0x1dbf848> collect = <built-in function collect> def f(*args, **kwargs): import pycuda.driver # appears to be idempotent, i.e. no harm in calling it more than once pycuda.driver.init() ctx = make_default_context() try: assert isinstance(ctx.get_device().name(), str) assert isinstance(ctx.get_device().compute_capability(), tuple) assert isinstance(ctx.get_device().get_attributes(), dict) > inner_f(*args, **kwargs) /usr/lib/python2.7/dist-packages/pycuda/tools.py:434: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = <test_driver.TestDriver instance at 0x24e7d88> @mark_cuda_test def test_register_host_memory(self): if drv.get_version() < (4,): from py.test import skip skip("register_host_memory only exists on CUDA 4.0 and later") import sys if sys.platform == "darwin": from py.test import skip skip("register_host_memory is not supported on OS X") a = drv.aligned_empty((2**20,), np.float64, alignment=4096) > drv.register_host_memory(a) E LogicError: cuMemHostRegister failed: invalid value test_driver.py:559: LogicError ==================== 1 failed, 20 passed in 116.85 seconds ===================== This test fails both on ION (GeForce 9400M, CC 1.1) and GeForce 460 (CC 2.1). I've compiled PyCUDA with gcc 4.8, run with kernel 3.9 and drivers 304.88. Regards. -- Tomasz Rybak GPG/PGP key ID: 2AD5 9860 Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860 http://member.acm.org/~tomaszrybak

5 years, 11 months

AWS summit in Berlin

by Tomasz Rybak

Hello. I'll be in Berlin on 2014-05-15, attending AWS Summit. If anyone wants to meet there, or in Berlin around that time, I'll gladly do so. Best regards. -- Tomasz Rybak GPG/PGP key ID: 2AD5 9860 Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860 http://member.acm.org/~tomaszrybak

6 years, 3 months

Can I use cuPrintf in PyCuda?

by 金陆

I am using an old PC with an old GPU card(GeForce 9800 GT). As you may know, 9800 does not support "printf" function in device code. However, Nvidia supplies cuPrintf. In CU file, it can be used like following. The situation is how can I use "cudaPrintfInit, cudaPrintfDisplay and cudaPrintfEnd" in Python code? Thanks [cu starts] #include "cuPrintf.cu" __global__ void my_kernel() { cuPrintf("x=%u, y=%u, offset=%u\t", x, y, offset); } int main() { cudaPrintfInit(); my_kernel<<<...>>>(); cudaPrintfDisplay(stdout, true); cudaPrintfEnd(); } [cu ends]

6 years, 3 months

Way to debug PyCuda code on Windows os?

by 金陆

I have read the FAQ of PyCUDA, but the bad news is that cuda-gdb does not run on Windows. So, my question is, how can I debug a PyCuda program on my Windows? Regards.

6 years, 3 months

C++ style comment yields bad code. Bug?

by 金陆

Hi, everyone There is a simple code, but I found if I use "//CUPRINTF("\n\n");" in test.py, the yielded kernel.cu file is invalid. That is so strange, because 1."//unsigned int y=threadIdx.y; " dose not ruin the yielded kernel.cu; 2. if I use "/*CUPRINTF("\n\n");*/", there is no problem I am using pycuda-2013.1.1.win32-py2.7.exe from http://www.lfd.uci.edu/~gohlke/pythonlibs/ on Winxp Thanks in advance [test.py starts] #coding=utf-8 import pycuda.driver as drv import pycuda.autoinit from pycuda.compiler import SourceModule def compile_kernel(kernel_code, kernel_name): mod = SourceModule(kernel_code) func = mod.get_function(kernel_name) return func my_kernel_code = """ #include "math.h" #include "stdio.h" __global__ void my_kernel() { unsigned int x=threadIdx.x; //unsigned int y=threadIdx.y; //CUPRINTF("\n\n"); } """ kernel = compile_kernel(my_kernel_code, 'my_kernel') if __name__ == '__main__': pass [test.py ends] [kernel.cu starts] extern "C" { #include "math.h" #include "stdio.h" __global__ void my_kernel() { unsigned int x=threadIdx.x; //unsigned int y=threadIdx.y; //CUPRINTF(" "); } } [kernel.cu ends]

6 years, 3 months

Re: [PyCUDA] RFC - a PEP for adding a dedicated matrix multiplication operator to Python

by Nathaniel Smith

Hi again, As an update for anyone not following python-ideas or numpy-discussion: Guido has accepted the addition of @ as a matrix multiplication operator in principle, once some details are worked out: https://mail.python.org/pipermail/python-ideas/2014-March/027109.html So if you have any feedback you've been sitting on, now is a good time to speak up :-). The current PEP draft can be seen here: http://legacy.python.org/dev/peps/pep-0465/ I've also started threads on numpy-discussion specifically about the details to be worked out, in case anyone wants to comment: Precedence/associativity of @: http://mail.scipy.org/pipermail/numpy-discussion/2014-March/069444.html Do we care about @@: http://mail.scipy.org/pipermail/numpy-discussion/2014-March/069448.html I believe that if you aren't subscribed to numpy-discussion, then you can still post through gmane (though unfortunately their anti-spam munging of the @ character makes these threads kind of hard to read!): http://news.gmane.org/gmane.comp.python.numeric.general Or worst case, just send me an email directly, and I'll collate. Cheers, -n On Sun, Mar 9, 2014 at 3:42 AM, Nathaniel Smith <njs(a)pobox.com> wrote: > Hey all, > > I just wanted to give you a heads-up about a draft Python Enhancement > Proposal that I'm working on, for adding a dedicated matrix > multiplication operator ('@') to Python: > https://github.com/numpy/numpy/pull/4351 > https://github.com/njsmith/numpy/blob/matmul-pep/doc/neps/return-of-revenge… > > This seems relevant to PyCUDA because IIUC your GPUArray is an ndarray > work-a-like, so if this PEP gets accepted then you'd probably want to > be implementing support for @. > > So, I'd definitely welcome feedback on any aspect of the document, but > especially the "Semantics" section (since whatever ends up there is > what numpy and other projects will be doing, and perhaps PyCUDA as > well?). And if you *really* like it, then please let me know, because > I'd love to be able to include PyCUDA on the list of projects > "planning to implement @", as further evidence of our awesome > consensus and how the BDFL should totally accept this PEP. Let me > know! > > -n > > -- > Nathaniel J. Smith > Postdoctoral researcher - Informatics - University of Edinburgh > http://vorpus.org -- Nathaniel J. Smith Postdoctoral researcher - Informatics - University of Edinburgh http://vorpus.org

6 years, 3 months

RFC - a PEP for adding a dedicated matrix multiplication operator to Python

by Nathaniel Smith

Hey all, I just wanted to give you a heads-up about a draft Python Enhancement Proposal that I'm working on, for adding a dedicated matrix multiplication operator ('@') to Python: https://github.com/numpy/numpy/pull/4351 https://github.com/njsmith/numpy/blob/matmul-pep/doc/neps/return-of-revenge… This seems relevant to PyCUDA because IIUC your GPUArray is an ndarray work-a-like, so if this PEP gets accepted then you'd probably want to be implementing support for @. So, I'd definitely welcome feedback on any aspect of the document, but especially the "Semantics" section (since whatever ends up there is what numpy and other projects will be doing, and perhaps PyCUDA as well?). And if you *really* like it, then please let me know, because I'd love to be able to include PyCUDA on the list of projects "planning to implement @", as further evidence of our awesome consensus and how the BDFL should totally accept this PEP. Let me know! -n -- Nathaniel J. Smith Postdoctoral researcher - Informatics - University of Edinburgh http://vorpus.org

6 years, 4 months

(no subject)

by 金陆

Hi, all the following code is translated from the book "CUDA by example". When it runs, I get 'pycuda._driver.LogicError: cuFuncSetBlockShape failed: invalid value' I have searched via google, and some pages said the the block is too big. But I only use 64, which is very small! Any helps? thanks from scipy import * from scipy import misc import pycuda.autoinit import pycuda.driver as drv from pycuda.compiler import SourceModule DIM = 64 mod = SourceModule(""" #define DIM 64 struct cuComplex{ float r, i; __device__ cuComplex(float a, float b):r(a),i(b){} __device__ float magnitude2() { return r*r+i*i;} __device__ cuComplex operator*(const cuComplex &a) { return cuComplex(r*a.r-i*a.i, i*a.r+r*a.i); } __device__ cuComplex operator+(const cuComplex &a) { return cuComplex(r + a.r, i + a.i); } }; __device__ int julia(int x, int y) { const float scale=1.5; float jx=scale*(float)(DIM/2-x)/(DIM/2); float jy=scale*(float)(DIM/2-y)/(DIM/2); cuComplex c(-0.8, 0.156), a(jx, jy); for(int i=0; i<200; i++) { a=a*a+c; if (a.magnitude2()>1000) return 0; } return 1; } __global__ void kernel(unsigned char *ptr) { unsigned int x=blockIdx.x; unsigned int y=blockIdx.y; unsigned int offset = x+y*gridDim.x; int juliaValue=julia(x, y); ptr[offset]=255*juliaValue; } """) mb = mod.get_function("kernel") dev_bitmap=zeros((DIM, DIM), dtype=dtype('c') ) mb( drv.InOut(dev_bitmap), block=(DIM, DIM, 1), grid = (1, 1)) misc.imsave("mandelbrot.png", dev_bitmap.T)

6 years, 4 months

Re: [PyCUDA] PyCuda - Bindless Textures

by Andreas Kloeckner

Dear John, John Holdener <jrholden(a)ucsc.edu> writes: > I'm working for a visualization lab at University of California Santa > Cruz, namely run by the High Performance AstroComputing Center. We > are using ray casting on a lot of AMR data, and would like to use > bindless textures. Are these supported by PyCuda? I see no questions >p or documentation on it. Thus far, there is no support for texture objects, but I'd be happy to take a patch. Andreas

6 years, 4 months

2020

2019

2018

2017

2016

2015

2014

2013

2012

2011

2010

2009

2008

PyCUDA March 2014