Bug in register_host_memory
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
4 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
5 years, 4 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]
5 years, 4 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.
5 years, 4 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]
5 years, 4 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-reven...
>
> 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
5 years, 4 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-reven...
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
5 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)
5 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
5 years, 4 months
