Converted NVIDIA RadixSort to PyOpenCL
by Ian Johnson
Hello all,
I finally bit the bullet and got radix working in PyOpenCL :)
It's also improved over the SDK example because it does keys and values,
mostly thanks to my advisor.
Additionally this sort will handle any size array as long as it is a power
of 2. The shipped example does not allow for arrays smaller than 32768, but
I've hooked up their naive scan to allow all smaller arrays.
https://github.com/enjalot/adventures_in_opencl/tree/master/experiments/r...
all you really need are radix.py, RadixSort.cl and Scan_b.cl
some simple tests are at the bottom of radix.py
I hammered this out because I need it for a project, it's not all that clean
and I didn't add support for sorting on keys only (altho it wouldn't take
much to add that, and I intend to at a later time when I need the
functionality). Hopefully this helps someone else out there. I'll also be
porting it using my own OpenCL C++ wrappers to include in my fluid
simulation library at some point.
I also began looking at AMD's radix from their SPH tutorial, but they use
local atomics which are not supported on my 9600M
--
Ian Johnson
http://enja.org
8 years
test_array results on OSX
by Randy Heiland
Hello,
Fwiw, I posted the outputs from the 'test_array.py' script at http://mypage.iu.edu/~heiland/pyopencl/
I realize the failures (and fixes) that occur here may be due to Andreas's lack of having an OSX platform to test on, so I don't intend for this post to be a nuisance. I'm simply curious what people are getting as results (from tests & examples) on various platforms and thought I'd offer up some from OSX.
Btw, if there is some guidance on debugging pyopencl, I'd certainly welcome it and try to contribute.
-Randy
8 years, 2 months
invalid work group size
by Randy Heiland
Hi all,
I'm hoping someone will offer some advice for a newbie. I'm getting the following error:
python demo_mandelbrot.py
Traceback (most recent call last):
File "demo_mandelbrot.py", line 161, in <module>
test = Mandelbrot()
File "demo_mandelbrot.py", line 120, in __init__
self.create_image()
File "demo_mandelbrot.py", line 147, in create_image
self.draw(-2.13, 0.77, -1.3, 1.3)
File "demo_mandelbrot.py", line 133, in draw
output = calc_fractal(q, maxiter)
File "demo_mandelbrot.py", line 66, in calc_fractal_opencl
output_opencl, np.uint16(maxiter))
File "/Library/Python/2.7/site-packages/pyopencl-2011.1.2-py2.7-macosx-10.7-intel.egg/pyopencl/__init__.py", line 221, in kernel_call
global_offset, wait_for, g_times_l=g_times_l)
pyopencl.LogicError: clEnqueueNDRangeKernel failed: invalid work group size
To avoid cluttering this email, I've provided a dump of my properties (AMD card) here:
http://mypage.iu.edu/~heiland/pyopencl/
I'm working out of git:
In [11]: pyopencl.version.VERSION
Out[11]: (2011, 1, 2)
and also have:
PyOpenGL-3.0.1-py2.7
I did a search here http://search.gmane.org/?query=invalid+work+group+size&group=gmane.comp.p...
but it didn't shed much light, for me anyway.
thanks, Randy
8 years, 4 months
Intel sdk verbosity
by Vincent Favre-Nicolin
Hi,
I've been using PyCUDA a lot (see http://pynx.sf.net) and I am just
switching to PyOpenCL, which seems to work very nicely, including with CPUs.
Right now I am working under Ubuntu with amd+intel+nvidia sdk's, which seem
to work nicely along each other and PyOpenCL.
Only slightly annoying thing, the intel compiler seems to be a bit verbose
when it does not need to be, e.g. I get messages like:
##########
Build on <pyopencl.Device 'Intel(R) Core(TM)2 Quad CPU Q9550 @ 2.83GHz' at
0x1af72c0> succeeded, but said:
Build started
Kernel <Fhkl> was successfully vectorized
Done.
##########
This is a pretty minor issue, but does someone know how to turn off this
message from the intel compiler (while still keeping real warnings, I don't
want to completely turn off PyOpenCL reporting of messages). I tried looking in
the intel manual, but did not find a compiler option for that.
Incidentally, the intel compiler seems much more efficient (3.5x speed) than
AMD's for my Intel Core 2 Q9550. Not sure why, either better vectorization or
faster sin/cos functions.
regards,
--
Vincent Favre-Nicolin http://inac.cea.fr
CEA/Grenoble Institut Nanosciences & Cryogénie
Laboratoire SP2M/Nano-structures et Rayonnement Synchrotron
17, rue des Martyrs
38054 Grenoble Cedex 9 - France
Université Joseph Fourier http://www.ujf-grenoble.fr
tél: (+33) 4 38 78 95 40 fax: (+33) 4 38 78 51 38
8 years, 4 months
Re: [PyOpenCL] Yapocis 0.2.3
by Andreas Kloeckner
On Mon, 29 Aug 2011 08:56:01 -0400, Sean True <sean.true(a)gmail.com> wrote:
> Has been uploaded to PyPI. It's still a bit raw, but adds a useful
> thin layer between Python and the PyOpenCL runtime.
>
> Declaring an interface:
>
> interface mandelbrot {
> kernel mandelbrot(in complex64 *q, outlike int16 *q, in int32 maxiter);
> }
>
> lets you call mandelbrot :
>
>
> output = calc_fractal(q, maxiter)
>
> No explicit buffer management, note that the output buffer q the
> return value, and that no declarations have to be made for it at all
> in the python code.
>
> I'm now using this actively for my personal work, and the package
> dependencies are primarily the same as PyOpenCL (pyparsing is the
> primary addition).
>
> Licensed PSF, use as you like.
This looks like good stuff. A friend of mine had a need for something
like this recently, and I was happy I could point him towards it. I've
also added a link to the FAQ.
Quick request: Please add a link to the PyPI page to the doc front page,
so that there's a central place to link to.
Andreas
8 years, 5 months
Yapocis 0.2.3
by Sean True
Has been uploaded to PyPI. It's still a bit raw, but adds a useful
thin layer between Python and the PyOpenCL runtime.
Declaring an interface:
interface mandelbrot {
kernel mandelbrot(in complex64 *q, outlike int16 *q, in int32 maxiter);
}
lets you call mandelbrot :
output = calc_fractal(q, maxiter)
No explicit buffer management, note that the output buffer q the
return value, and that no declarations have to be made for it at all
in the python code.
I'm now using this actively for my personal work, and the package
dependencies are primarily the same as PyOpenCL (pyparsing is the
primary addition).
Licensed PSF, use as you like.
-- Sean
Sean True
Swapwizard Consulting
8 years, 5 months
PyCUDA, PyOpenCL in Debian, Ubuntu
by Tomasz Rybak
Hello.
I am sending one email to two groups - just call me lazy ;-)
Yesterday python-pycuda got accepted into Debian unstable,
and today it landed in archives.
This means that Debian unstable contains PyCUDA 2011.1.3 with
git patches from 2011-08-13 (smem_alloc_granularity) and
from 2011-08-14 (fix of debug code leak) is available in Debian
unstable. Package should migrate into testing (and thus
become candidate for inclusion into next stable Debian)
after two weeks if there are no problems with it.
At the same time both Debian unstable and Debian testing
both contain PyOpenCL 2011.1.2 with changes from git
up to 2011-07-16. After 2011.2 is released I intend to
upload it into Debian.
Unfortunately Ubuntu still contains PyOpenCL 0.92.
Automatic upload of Debian version into Ubuntu did
not worked, probably because of bug
https://bugs.launchpad.net/ubuntu/+source/pyopencl/+bug/763457
which forces users to install NVIDIA drivers even
if they use another OpenCL provider. Debian version
contains fix (at least I hope that it fixes this problem)
but I got no reply to my message asking whether problem
is fixed or not.
As for PyCUDA in Ubuntu, it does not look like it will
be included soon. Ubuntu has different driver architecture,
and does not contain NVIDIA CUDA toolkit, and both are heavily
used by python-pycuda.
Best regards.
--
Tomasz Rybak <bogomips(a)post.pl> GPG/PGP key ID: 2AD5 9860
Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860
http://member.acm.org/~tomaszrybak
8 years, 5 months
test_array.py error (memory pool)
by Tomasz Rybak
Hello.
I have found another error.
I have been testing on ION (GeForce 9400M, CC 1.1, without double
support) and one of the tests (test_mem_pool_with_arrays) gives error:
============================= test session starts
==============================
platform linux2 -- Python 2.6.7 -- pytest-1.3.4
test path 1: test_array.py
test_array.py ...........................................................F
=================================== FAILURES
===================================
test_mem_pool_with_arrays[ctx_factory=<context factory for
<pyopencl.Device 'ION' at 0x2f38f80>>]
ctx_factory = <pyopencl.tools.ContextFactory instance at 0x311c050>
@pytools.test.mark_test.opencl
def test_mem_pool_with_arrays(ctx_factory):
context = ctx_factory()
queue = cl.CommandQueue(context)
mem_pool = cl_tools.MemoryPool(cl_tools.CLAllocator(context))
> a_dev = cl_array.arange(queue, 2000, dtype=np.float64,
allocator=mem_pool)
test_array.py:631:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
def arange(*args, **kwargs):
"""Create an array filled with numbers spaced `step` apart,
starting from `start` and ending at `stop`.
For floating point arguments, the length of the result is
`ceil((stop - start)/step)`. This rule may result in the last
element of the result being greater than stop.
"""
if isinstance(args[0], cl.Context):
from warnings import warn
warn("Passing a context as first argument is deprecated. "
"This will be continue to be accepted througout "
"versions 2011.x of PyOpenCL.",
DeprecationWarning, 2)
args = args[1:]
> return _arange(*args, **kwargs)
/usr/lib/python2.6/dist-packages/pyopencl/array.py:813:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
queue = <pyopencl._cl.CommandQueue object at 0x3125578>
def _arange(queue, *args, **kwargs):
# argument processing
-----------------------------------------------------
# Yuck. Thanks, numpy developers. ;)
from pytools import Record
class Info(Record):
pass
explicit_dtype = False
inf = Info()
inf.start = None
inf.stop = None
inf.step = None
inf.dtype = None
inf.allocator = None
if isinstance(args[-1], np.dtype):
dtype = args[-1]
args = args[:-1]
explicit_dtype = True
argc = len(args)
if argc == 0:
raise ValueError, "stop argument required"
elif argc == 1:
inf.stop = args[0]
elif argc == 2:
inf.start = args[0]
inf.stop = args[1]
elif argc == 3:
inf.start = args[0]
inf.stop = args[1]
inf.step = args[2]
else:
raise ValueError, "too many arguments"
admissible_names = ["start", "stop", "step", "dtype",
"allocator"]
for k, v in kwargs.iteritems():
if k in admissible_names:
if getattr(inf, k) is None:
setattr(inf, k, v)
if k == "dtype":
explicit_dtype = True
else:
raise ValueError, "may not specify '%s' by position
and keyword" % k
else:
raise ValueError, "unexpected keyword argument '%s'" % k
if inf.start is None:
inf.start = 0
if inf.step is None:
inf.step = 1
if inf.dtype is None:
inf.dtype = np.array([inf.start, inf.stop, inf.step]).dtype
# actual functionality
----------------------------------------------------
dtype = np.dtype(inf.dtype)
start = dtype.type(inf.start)
step = dtype.type(inf.step)
stop = dtype.type(inf.stop)
if not explicit_dtype:
raise TypeError("arange requires dtype argument")
from math import ceil
size = int(ceil((stop-start)/step))
result = Array(queue, (size,), dtype, allocator=inf.allocator)
> _arange_knl(result, start, step, queue=queue)
/usr/lib/python2.6/dist-packages/pyopencl/array.py:790:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
def kernel_runner(*args, **kwargs):
repr_ary = args[0]
queue = kwargs.pop("queue", None) or repr_ary.queue
gs, ls = repr_ary.get_sizes(queue)
> knl = kernel_getter(*args)
/usr/lib/python2.6/dist-packages/pyopencl/array.py:147:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
result = array([ 0.00000000e+000, 5.26354425e-315, 0.00000000e
+000, ...,
1.76537186e-304, 1.76537882e-304, 1.76538751e-304])
start = 0.0, step = 1.0
@elwise_kernel_runner
def _arange_knl(result, start, step):
return elementwise.get_arange_kernel(
> result.context, result.dtype)
/usr/lib/python2.6/dist-packages/pyopencl/array.py:719:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
context = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
dtype = dtype('float64')
> ???
<string>:2:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
func = <function get_arange_kernel at 0x2dc8d70>
context = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
@decorator
def first_arg_dependent_memoize(func, context, *args):
"""Provides memoization for things that get created inside
a context, i.e. mainly programs and kernels. Assumes that
the first argument of the decorated function is the context.
"""
try:
ctx_dict = func._pyopencl_first_arg_dep_memoize_dic
except AttributeError:
# FIXME: This may keep contexts alive longer than desired.
# But I guess since the memory in them is freed, who cares.
ctx_dict = func._pyopencl_first_arg_dep_memoize_dic = {}
try:
return ctx_dict[context][args]
except KeyError:
first_arg_dependent_memoized_functions.append(func)
arg_dict = ctx_dict.setdefault(context, {})
> result = func(context, *args)
/usr/lib/python2.6/dist-packages/pyopencl/tools.py:69:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
context = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
dtype = dtype('float64')
@context_dependent_memoize
def get_arange_kernel(context, dtype):
return get_elwise_kernel(context,
"%(tp)s *z, %(tp)s start, %(tp)s step" % {
"tp": dtype_to_ctype(dtype),
},
"z[i] = start + i*step",
> name="arange")
/usr/lib/python2.6/dist-packages/pyopencl/elementwise.py:391:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
context = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
arguments = 'double *z, double start, double step'
operation = 'z[i] = start + i*step', name = 'arange', options = []
def get_elwise_kernel(context, arguments, operation,
name="elwise_kernel", options=[], **kwargs):
"""Return a L{pyopencl.Kernel} that performs the same scalar
operation
on one or several vectors.
"""
func, arguments = get_elwise_kernel_and_types(
context, arguments, operation,
> name=name, options=options, **kwargs)
/usr/lib/python2.6/dist-packages/pyopencl/elementwise.py:113:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
context = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
arguments = 'double *z, double start, double step'
operation = 'z[i] = start + i*step', name = 'arange', options = []
preamble = '#pragma OPENCL EXTENSION cl_khr_fp64: enable\n\n\n'
def get_elwise_kernel_and_types(context, arguments, operation,
name="elwise_kernel", options=[], preamble="", **kwargs):
if isinstance(arguments, str):
from pyopencl.tools import parse_c_arg
parsed_args = [parse_c_arg(arg) for arg in
arguments.split(",")]
else:
parsed_args = arguments
for arg in parsed_args:
if np.float64 == arg.dtype:
preamble = (
"#pragma OPENCL EXTENSION cl_khr_fp64: enable\n
\n\n"
+ preamble)
break
parsed_args.append(ScalarArg(np.uintp, "n"))
prg = get_elwise_program(
context, parsed_args, operation,
> name=name, options=options, preamble=preamble, **kwargs)
/usr/lib/python2.6/dist-packages/pyopencl/elementwise.py:91:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
context = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
arguments = [<[TypeError("not enough arguments for format string")
raised in repr()] VectorArg object at 0x3363fc8>, <[TypeError("... at
0x3363170>, <[TypeError("not enough arguments for format string") raised
in repr()] ScalarArg object at 0x3363050>]
operation = 'z[i] = start + i*step', name = 'arange', options = []
preamble = '#pragma OPENCL EXTENSION cl_khr_fp64: enable\n\n\n',
loop_prep = ''
after_loop = ''
def get_elwise_program(context, arguments, operation,
name="elwise_kernel", options=[],
preamble="", loop_prep="", after_loop=""):
from pyopencl import Program
source = ("""
%(preamble)s
__kernel void %(name)s(%(arguments)s)
{
unsigned lid = get_local_id(0);
unsigned gsize = get_global_size(0);
unsigned work_item_start =
get_local_size(0)*get_group_id(0);
unsigned i;
%(loop_prep)s;
for (i = work_item_start + lid; i < n; i += gsize)
{
%(operation)s;
}
%(after_loop)s;
}
""" % {
"arguments": ", ".join(arg.declarator() for arg in
arguments),
"operation": operation,
"name": name,
"preamble": preamble,
"loop_prep": loop_prep,
"after_loop": after_loop,
})
> return Program(context, source).build(options)
/usr/lib/python2.6/dist-packages/pyopencl/elementwise.py:69:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
self = <pyopencl.Program object at 0x349be10>
options = ['-I',
'/usr/lib/python2.6/dist-packages/pyopencl/../../../../include/pyopencl']
devices = None, cache_dir = None
def build(self, options=[], devices=None, cache_dir=None):
if isinstance(options, str):
options = [options]
options = options + ["-I", _find_pyopencl_include_path()]
if self._prg is not None:
self._prg._build(options, devices)
else:
from pyopencl.cache import
create_built_program_from_source_cached
self._prg = create_built_program_from_source_cached(
self._context, self._source, options, devices,
> cache_dir=cache_dir)
/usr/lib/python2.6/dist-packages/pyopencl/__init__.py:462:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
ctx = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
src = '\n #pragma OPENCL EXTENSION cl_khr_fp64: enable\n\n\n\n\n
__kernel void arange(__global double *z, doub...; i < n; i += gsize)\n
{\n z[i] = start + i*step;\n }\n\n ;
\n }\n '
options = ['-I',
'/usr/lib/python2.6/dist-packages/pyopencl/../../../../include/pyopencl']
devices = None, cache_dir = None
def create_built_program_from_source_cached(ctx, src, options=[],
devices=None,
cache_dir=None):
try:
if cache_dir != False:
prg, already_built =
_create_built_program_from_source_cached(
> ctx, src, options, devices, cache_dir)
/usr/lib/python2.6/dist-packages/pyopencl/cache.py:436:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
ctx = <pyopencl.Context at 0x339d3b0 on <pyopencl.Device 'ION' at
0x2f38f80>>
src = '\n #pragma OPENCL EXTENSION cl_khr_fp64: enable\n\n\n\n\n
__kernel void arange(__global double *z, doub...; i < n; i += gsize)\n
{\n z[i] = start + i*step;\n }\n\n ;
\n }\n '
options = ['-I',
'/usr/lib/python2.6/dist-packages/pyopencl/../../../../include/pyopencl']
devices = [<pyopencl.Device 'ION' at 0x2f38f80>]
cache_dir = '/tmp/pyopencl-compiler-cache-v2-uidtomus-py2.6.7.final.0'
def _create_built_program_from_source_cached(ctx, src, options,
devices, cache_dir):
include_path = ["."] + [
option[2:]
for option in options
if option.startswith("-I") or option.startswith("/I")]
if cache_dir is None:
from os.path import join
from tempfile import gettempdir
import getpass
cache_dir = join(gettempdir(),
"pyopencl-compiler-cache-v2-uid%s-py%s" % (
getpass.getuser(), ".".join(str(i) for i in
sys.version_info)))
# {{{ ensure cache directory exists
try:
os.mkdir(cache_dir)
except OSError, e:
from errno import EEXIST
if e.errno != EEXIST:
raise
# }}}
if devices is None:
devices = ctx.devices
cache_keys = [get_cache_key(device, options, src) for device in
devices]
binaries = []
to_be_built_indices = []
logs = []
for i, (device, cache_key) in enumerate(zip(devices,
cache_keys)):
cache_result = retrieve_from_cache(cache_dir, cache_key)
if cache_result is None:
to_be_built_indices.append(i)
binaries.append(None)
logs.append(None)
else:
binary, log = cache_result
binaries.append(binary)
logs.append(log)
message = (75*"="+"\n").join(
"Build on %s succeeded, but said:\n\n%s" % (dev, log)
for dev, log in zip(devices, logs)
if log is not None and log.strip())
if message:
from warnings import warn
warn("Build succeeded, but resulted in non-empty logs:
\n"+message)
# {{{ build on the build-needing devices, in one go
result = None
already_built = False
if to_be_built_indices:
prg = _cl._Program(ctx, src)
> prg.build(options, [devices[i] for i in
to_be_built_indices])
/usr/lib/python2.6/dist-packages/pyopencl/cache.py:361:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
_ _ _ _
self = <pyopencl._cl._Program object at 0x7ff712cc8520>
options =
'-I /usr/lib/python2.6/dist-packages/pyopencl/../../../../include/pyopencl'
devices = [<pyopencl.Device 'ION' at 0x2f38f80>]
def program_build(self, options=[], devices=None):
if isinstance(options, list):
options = " ".join(options)
err = None
try:
self._build(options=options, devices=devices)
except Exception, e:
from pytools import Record
class ErrorRecord(Record):
pass
what = e.what + "\n\n" + (75*"="+"\n").join(
"Build on %s:\n\n%s" % (dev, log)
for dev, log in self._get_build_logs())
code = e.code
routine = e.routine
err = _cl.RuntimeError(
ErrorRecord(
what=lambda : what,
code=lambda : code,
routine=lambda : routine))
if err is not None:
# Python 3.2 outputs the whole tree of currently active
exceptions
# This serves to remove one (redundant) level from that
nesting.
> raise err
E RuntimeError: clBuildProgram failed: invalid binary -
E
E Build on <pyopencl.Device 'ION' at 0x2f38f80>:
E
E ptxas application ptx input, line 52; error : Instruction
'cvt' requires SM 1.3 or higher, or map_f64_to_f32 directive
E ptxas application ptx input, line 53; error : Instruction
'ld' requires SM 1.3 or higher, or map_f64_to_f32 directive
E ptxas application ptx input, line 54; error : Instruction
'mul' requires SM 1.3 or higher, or map_f64_to_f32 directive
E ptxas application ptx input, line 55; error : Instruction
'ld' requires SM 1.3 or higher, or map_f64_to_f32 directive
E ptxas application ptx input, line 56; error : Instruction
'add' requires SM 1.3 or higher, or map_f64_to_f32 directive
E ptxas application ptx input, line 60; error : Instruction
'st' requires SM 1.3 or higher, or map_f64_to_f32 directive
E ptxas fatal : Ptx assembly aborted due to errors
E ptxas application ptx input, line 52; warning : Double is
not supported. Demoting to float
/usr/lib/python2.6/dist-packages/pyopencl/__init__.py:145: RuntimeError
===================== 1 failed, 59 passed in 97.97 seconds
=====================
After changing line 631 from
a_dev = cl_array.arange(queue, 2000, dtype=np.float64,
allocator=mem_pool)
to
a_dev = cl_array.arange(queue, 2000, dtype=np.float32,
allocator=mem_pool)
(i.e. to dtype=np.float32) everything passes.
Best regards.
--
Tomasz Rybak <bogomips(a)post.pl> GPG/PGP key ID: 2AD5 9860
Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860
http://member.acm.org/~tomaszrybak
8 years, 5 months
Re: [PyOpenCL] PyOpenCL Digest, Vol 24, Issue 15 Re: error importing pyopencl
by M.Gelman
Jan,
Thanks for that tip. I realized that after reading the OpenCL specification
that type "double" is not required by the specification, so users should
check with their hardware documentation, both Nvidia and ATI, to determine
whether double is supported. In my case it is not.
Thanks, Max
On Tue, Aug 23, 2011 at 3:00 PM, <pyopencl-request(a)tiker.net> wrote:
> Send PyOpenCL mailing list submissions to
> pyopencl(a)tiker.net
>
> To subscribe or unsubscribe via the World Wide Web, visit
> http://lists.tiker.net/listinfo/pyopencl
> or, via email, send a message with subject or body 'help' to
> pyopencl-request(a)tiker.net
>
> You can reach the person managing the list at
> pyopencl-owner(a)tiker.net
>
> When replying, please edit your Subject line so it is more specific
> than "Re: Contents of PyOpenCL digest..."
>
>
> Today's Topics:
>
> 1. Re: yapocis? Yet another python opencl interface
> specification layer? (Sean True)
> 2. Re: error importing pyopencl (Jan Meinke)
>
>
> ----------------------------------------------------------------------
>
> Message: 1
> Date: Mon, 22 Aug 2011 20:56:51 +0000 (UTC)
> From: Sean True <sean.true(a)gmail.com>
> To: pyopencl(a)tiker.net
> Subject: Re: [PyOpenCL] yapocis? Yet another python opencl interface
> specification layer?
> Message-ID: <loom.20110822T225431-979(a)post.gmane.org>
> Content-Type: text/plain; charset=us-ascii
>
> Andreas:
>
> http://yapocis.readthedocs.org/ and https://github.com/seantrue/Yapocishave rev 0.1 of the layer I
> mentioned before.
>
> I'd be interested in your reaction ;-)
>
> -- Sean
>
>
>
>
>
>
> ------------------------------
>
> Message: 2
> Date: Mon, 22 Aug 2011 22:52:28 +0200
> From: Jan Meinke <jan.meinke(a)gmail.com>
> To: "M.Gelman" <m.gelman08(a)gmail.com>
> Cc: pyopencl(a)tiker.net
> Subject: Re: [PyOpenCL] error importing pyopencl
> Message-ID:
> <CAEFwH5DGD8CzdTjzLxLibNt7-_355JMOiH=rSi_=T8+Li-8Zdw(a)mail.gmail.com
> >
> Content-Type: text/plain; charset="iso-8859-1"
>
> Dear Max,
>
> AMD does not enable double precision support in OpenCL kernels by default.
> Instead you need to add
>
> #pragma OPENCL EXTENSION cl_amd_fp64 : enable
>
> to your kernel. This should work on the CPU but it will break the
> compilation on your GPU since it doesn't support double precision as far as
> I know.
>
> Jan
>
>
> On Thu, Aug 4, 2011 at 1:19 AM, M.Gelman <m.gelman08(a)gmail.com> wrote:
>
> > Thanks for your suggestion. I was able to fix it by copying the
> > libOpenCL.so provided by the ati stream sdk to /usr/lib.
> >
> > However, running the benchmark-all.py. I get the following errors about a
> > "double" type not enabled. It looks to me like everything is working
> > otherwise, including the other tests.
> >
> > ('Execution time of test without OpenCL: ', 8.684561014175415, 's')
> > ===============================================================
> > ('Platform name:', 'ATI Stream')
> > ('Platform profile:', 'FULL_PROFILE')
> > ('Platform vendor:', 'Advanced Micro Devices, Inc.')
> > ('Platform version:', 'OpenCL 1.1 ATI-Stream-v2.3 (451)')
> > ---------------------------------------------------------------
> > ('Device name:', 'ATI RV710')
> > ('Device type:', 'GPU')
> > ('Device memory: ', 256, 'MB')
> > ('Device max clock speed:', 600, 'MHz')
> > ('Device compute units:', 2)
> >
> /usr/local/lib/python2.6/dist-packages/pyopencl-2011.1.2-py2.6-linux-x86_64.egg/pyopencl/cache.py:343:
> > UserWarning: Build succeeded, but resulted in non-empty logs:
> > Build on <pyopencl.Device 'ATI RV710' at 0x21a9b20> succeeded, but said:
> >
> > /tmp/OCLQdtqga.cl(11): warning: double-precision constant is represented
> as
> > single-precision constant because double is not enabled
> > c[gid] = c[gid] * (a[gid] / 2.0);
> > ^
> >
> >
> > warn("Build succeeded, but resulted in non-empty logs:\n"+message)
> > Execution time of test: 0.00788741 s
> > benchmark-all.py:70: DeprecationWarning: 'enqueue_read_buffer' has been
> > deprecated in version 2011.1. Please use enqueue_copy() instead.
> > cl.enqueue_read_buffer(queue, dest_buf, c).wait()
> > Results OK
> > ===============================================================
> > ('Platform name:', 'ATI Stream')
> > ('Platform profile:', 'FULL_PROFILE')
> > ('Platform vendor:', 'Advanced Micro Devices, Inc.')
> > ('Platform version:', 'OpenCL 1.1 ATI-Stream-v2.3 (451)')
> > ---------------------------------------------------------------
> > ('Device name:', 'Intel(R) Core(TM) i7 CPU 870 @ 2.93GHz')
> > ('Device type:', 'CPU')
> > ('Device memory: ', 3072, 'MB')
> > ('Device max clock speed:', 2934, 'MHz')
> > ('Device compute units:', 8)
> >
> /usr/local/lib/python2.6/dist-packages/pyopencl-2011.1.2-py2.6-linux-x86_64.egg/pyopencl/cache.py:343:
> > UserWarning: Build succeeded, but resulted in non-empty logs:
> > Build on <pyopencl.Device 'Intel(R) Core(TM) i7 CPU 870 @
> 2.93GHz'
> > at 0x21bb140> succeeded, but said:
> >
> > /tmp/OCLLl8sG4.cl(11): warning: double-precision constant is represented
> as
> > single-precision constant because double is not enabled
> > c[gid] = c[gid] * (a[gid] / 2.0);
> > ^
> >
> >
> > warn("Build succeeded, but resulted in non-empty logs:\n"+message)
> > Execution time of test: 0.00134482 s
> > Results OK
> >
> >
> > On Wed, Aug 3, 2011 at 4:36 AM, Andreas Kloeckner <
> lists(a)informa.tiker.net
> > > wrote:
> >
> >> On Tue, 2 Aug 2011 19:05:33 -0400, "M.Gelman" <m.gelman08(a)gmail.com>
> >> wrote:
> >> Non-text part: multipart/alternative
> >> > I have been trying to get pyopencl to work with my ati card and ati
> >> stream
> >> > for awhile. I followed the tutorial an was able to get a good
> >> compilation
> >> > using this script:
> >> >
> >> > python configure.py \
> >> > --boost-inc-dir=/usr/include/boost \
> >> > --boost-lib-dir=/usr/lib \
> >> > --boost-python-libname=boost_python-mt-py26 \
> >> > --cl-inc-dir=/opt/ati-stream-sdk-v2.3-lnx64/include \
> >> > --cl-lib-dir=/opt/ati-stream-sdk-v2.3-lnx64/lib/x86_64 \
> >> > --cl-libname=OpenCL
> >> >
> >> > However when I import pyopencl, I get:
> >> >
> >> > >>> import pyopencl
> >> > Traceback (most recent call last):
> >> > File "<stdin>", line 1, in <module>
> >> > File
> >> >
> >>
> "/usr/local/lib/python2.6/dist-packages/pyopencl-2011.1.2-py2.6-linux-x86_64.egg/pyopencl/__init__.py",
> >> > line 4, in <module>
> >> > import pyopencl._cl as _cl
> >> > ImportError:
> >> >
> >>
> /usr/local/lib/python2.6/dist-packages/pyopencl-2011.1.2-py2.6-linux-x86_64.egg/pyopencl/_cl.so:
> >> > symbol clEnqueueWriteBufferRect, version OPENCL_1.1 not defined in
> file
> >> > libOpenCL.so with link time reference
> >> >
> >> > Its something probably simple however, I cannot get passed it.
> >> >
> >> > Here are my env vars:
> >> > export ATISTREAMSDKROOT=/opt/ati-stream-sdk-v2.3-lnx64
> >> > export ATISTREAMSDKSAMPLEROOT=/opt/ati-stream-sdk-v2.3-lnx64
> >> > export LD_LIBRARY_PATH=$ATISTREAMSDKROOT/lib/x86_64:$LD_LIBRARY_PATH
> >>
> >> Your OpenCL header does not match your OpenCL library. Check with
> >>
> >> $ ldd
> >>
> >>
> /usr/local/lib/python2.6/dist-packages/pyopencl-2011.1.2-py2.6-linux-x86_64.egg/pyopencl/_cl.so
> >> (one line)
> >>
> >> to see whether the library is the one you expect.
> >>
> >> HTH,
> >> Andreas
> >>
> >> PS: Please send email to pyopencl@, not pyopencl-owner@.
> >>
> >
> >
> >
> > --
> > M.Gelman
> > (412)-540-5238
> > LinkedIn Profile <http://www.linkedin.com/in/maxgelman>
> >
> > _______________________________________________
> > PyOpenCL mailing list
> > PyOpenCL(a)tiker.net
> > http://lists.tiker.net/listinfo/pyopencl
> >
> >
>
8 years, 5 months
yapocis? Yet another python opencl interface specification layer?
by Sean True
I'm wondering if yet-another-python-opencl-interface-layer would be of interest.
This layer is built on top of pyopencl, and is intended to use familiar
RPC definitions to ease the special joys of talking to opencl.
I've implemented this on OS/X and it is working for me and has been stable, oh,
for several days ;-) It eliminates the need for routine glue code,
and can be extended to minimize movement of data between host and opencl service.
This code is currently private, and I will be happy to either fork pyopencl and
submit a pull after integration, or a separate github project.
-- Sean
Sean True
Swapwizard Consulting
Presumed highlights:
# Use the rpc extension to define and load the kernel as a callable.
from rpc import kernels, interfaces
calc_fractal_opencl = kernels.loadProgram(interfaces.mandelbrot).mandelbrot
# Call it the way we like to call Python callables:
output = calc_fractal(q, maxiter)
# RPC definition language loosely based on Apollo NCS/OSF DCE/Microsoft IDL
# outlike is a novel keyword that says: allocate for me, return as part
of return vals.
interface mandelbrot {
kernel mandelbrot(in complex64 *q, outlike int16 *q, in int32 maxiter);
}
# mandelbrot.mako is just what it has always been.
__kernel void mandelbrot(__global float2 *q,
__global short *output, int const maxiter)
{
int gid = get_global_id(0);
float nreal, real = 0;
float imag = 0;
output[gid] = 0;
for(int curiter = 0; curiter < maxiter; curiter++) {
nreal = real*real - imag*imag + q[gid].x;
imag = 2* real*imag + q[gid].y;
real = nreal;
if (real*real + imag*imag > 4.0f)
output[gid] = curiter;
}
}
8 years, 5 months