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
7 years
How to provide precompiled PyOpenCL?
by Gonzo Kaczykowski
As in subject... I would like to install PyOpenCL on Windows machine which has no compiler. is there any possibility to build PyOpenCL on one machine and simply install on another?
Regards
7 years, 1 month
PyOpenCL and compyte - status?
by Tomasz Rybak
Hello.
I am writing with rather naive question, but what is
the status of those two, and when can we expect new
stable version?
I am asking this because I am trying to build Debian
package of PyOpenCL. Finally Debian has AMD/ATI OpenCL
library, so I want to change PyOpenCL package to take
this into consideration and not to require NVIDIA OpenCL
libraries. For now, PyOpenCL tests crash:
============================= test session starts
==============================
platform linux2 -- Python 2.7.2 -- pytest-1.3.4
test path 1: test_array.py
test_array.py E
==================================== ERRORS
====================================
ERROR collecting /usr/src/cuda/pyopencl/pyopencl-2011.1.2
+git20111121/test/test_array.py
test_array.py:669: in <module>
> register_dtype(mmc_dtype, "minmax_collector")
/usr/lib/python2.7/dist-packages/pyopencl/compyte/dtypes.py:50: in
register_dtype
> % (dtype, DTYPE_TO_NAME[dtype], ", ".join(c_names)))
E RuntimeError: dtype '[('cur_min', '<i4'), ('cur_max',
'<i4'), ('pad', '<i4')]' already registered (as 'minmax_collector', new
names 'minmax_collector')
=========================== 1 error in 0.21 seconds
============================
So my question is - should I just wait for few days till situation
gets better, or is this more serious, requiring more time to fix,
and I should just package older version for now?
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
7 years, 1 month
pyopencl code completion in eclipse pydev
by Dieter Morgenroth
Hi,
Did anyone get the pydev code completion to work with pyopencl? Adding
pyopencl to the "Forced Built ins" list didnt help. I dont get any code
completion and eclipse shows me a lot of errors because it doesnt find
the inputs.
Or is there a trick to get at least rid of the error messages?
Best regards,
Dieter
7 years, 1 month
Re: [PyOpenCL] Ghosting with bilinear Interpolation of OpenCV-Images
by Andreas Tharang
It wasn't a driver bug.
It has something to do with the way OpenCV stores the image in his
array-structure.
The following code should be the solution for this thread:
import pyopencl as cl
import numpy as np
import cv2 # OpenCV 2.3.1
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
clImageFormat = cl.ImageFormat(cl.channel_order.R,
cl.channel_type.UNORM_INT8)
In_host = cv2.cvtColor(cv2.imread("Test.jpg"), cv2.COLOR_BGR2GRAY)
aSize = In_host.shape[1::-1]
In_buf = cl.Image(ctx, cl.mem_flags.READ_ONLY |
cl.mem_flags.COPY_HOST_PTR, clImageFormat, aSize, None, In_host.tostring())
Out_host = np.empty(aSize, np.uint8)
Out_buf = cl.Image(ctx, cl.mem_flags.WRITE_ONLY, clImageFormat, aSize)
prg = cl.Program(ctx, """
const sampler_t smp = CLK_NORMALIZED_COORDS_TRUE |
CLK_FILTER_LINEAR | CLK_ADDRESS_NONE;
__kernel void ImageCopy(__read_only image2d_t Img, __write_only
image2d_t Out)
{
const int2 dims = get_image_dim(Img);
const int2 Coords = (int2)(get_global_id(0), get_global_id(1));
const float2 NormCoords = (convert_float2(Coords) +
(float2)(1.7f, 0.3f)) / convert_float2(dims); // = (x + 1.7)/h ; (y + 0.3)/w
float4 Pixel = read_imagef(Img, smp, NormCoords);
write_imagef(Out, Coords, Pixel);
}""").build()
prg.ImageCopy(queue, aSize, None, In_buf, Out_buf)
cl.enqueue_read_image(queue, Out_buf, (0, 0, 0), aSize, Out_host).wait()
cv2.imwrite("out.png", np.reshape(Out_host, aSize[1::-1]))
thanks to all
7 years, 1 month
Re: [PyOpenCL] Loading binary does not work
by Andreas Kloeckner
On Tue, 22 Nov 2011 07:57:24 +0100, Matthias Vogelgesang <matthias.vogelgesang(a)gmail.com> wrote:
> Hi,
>
> 2011/11/22 Andreas Kloeckner <lists(a)informa.tiker.net>:
> >
> > Fixed in git. The problem was that options defaulted to a list, which
> > the backend code didn't know how to process.
>
> Much appreciated, thanks.
Btw--PyOpenCL already already caches the binaries that get compiled. So
if you're after time savings by building from binary, there likely
won't be any. Just FYI. :)
Andreas
7 years, 1 month
Re: [PyOpenCL] Ghosting with bilinear Interpolation of OpenCV-Images
by Andreas Tharang
At 01.11.2011 14:06, Frédéric Bastien wrote:
>
> The cuda develloper version of the driver always worked correctly when
> some others version failled for me. So you can try it if you want to
> check that powsobility. For the hardware, can you try on another computer?
>
> Fred
>
Now i've tried it on a ATI Radeon HD 6970. And ... - it works!
So it's really seems to be a driver or hardware bug.
The problem for me is that the developer-driver ( v.197.16) is not
compatible for my device and the new VERDE v.285.62-Driver freezes my
system after installation (on a fresh installed Win 7 x64 Pro).
This is why i have to use the v.280.26.
I've wrote this the NVidia-Support-Team but still got no answer.
Andreas
7 years, 1 month
Re: [PyOpenCL] Loading binary does not work
by Andreas Kloeckner
On Mon, 21 Nov 2011 21:09:20 +0100, Matthias Vogelgesang <matthias.vogelgesang(a)gmail.com> wrote:
> Hi folks,
>
> according to the reference, it should be possible to re-load the
> output from a built program. However, the following short code is not
> working
>
> ctx = cl.create_some_context(False)
> platform = cl.get_platforms()
> device = platform[0].get_devices()[0]
>
> program = cl.Program(ctx, """
> __kernel void simple(__global float *in, __global float *out)
> {
> out[get_global_id(0)] = in[get_global_id(0)];
> }""")
> program.build()
> binary = program.get_info(cl.program_info.BINARIES)[0]
>
> foo = cl.Program(ctx, [device], [binary])
> # I cannot build ...
> foo.build()
> # ... nor call the kernel
> foo.simple()
>
> Is there something wrong with this code or PyOpenCL?
Fixed in git. The problem was that options defaulted to a list, which
the backend code didn't know how to process.
Andreas
7 years, 1 month
Loading binary does not work
by Matthias Vogelgesang
Hi folks,
according to the reference, it should be possible to re-load the
output from a built program. However, the following short code is not
working
ctx = cl.create_some_context(False)
platform = cl.get_platforms()
device = platform[0].get_devices()[0]
program = cl.Program(ctx, """
__kernel void simple(__global float *in, __global float *out)
{
out[get_global_id(0)] = in[get_global_id(0)];
}""")
program.build()
binary = program.get_info(cl.program_info.BINARIES)[0]
foo = cl.Program(ctx, [device], [binary])
# I cannot build ...
foo.build()
# ... nor call the kernel
foo.simple()
Is there something wrong with this code or PyOpenCL?
Regards,
Matthias
--
Matthias Vogelgesang
Public-Key: http://tinyurl.com/2qcydl
7 years, 1 month
Re: [PyOpenCL] re. pyopencl against Intel
by Andreas Kloeckner
On Sun, 20 Nov 2011 05:07:55 -0800, gliptak <reply+m-20055400-1a725a8ae8e73dce2eb829a1d162d7a448aac846-352067(a)reply.github.com> wrote:
> I guess you didn't say it isn't supported, but it wasn't tested against.
>
> http://wiki.tiker.net/PyOpenCL/FrequentlyAskedQuestions#Which_implementat...
>
> In test/ test_clmath.py runs OK.
>
> Here are the fails I do see:
>
> $ LD_LIBRARY_PATH=/usr/lib64 python test_wrapper.py
> ============================= test session starts ==============================
> platform linux2 -- Python 2.7.2 -- pytest-2.2.0
> collected 12 items
>
> test_wrapper.py .F...s......
>
> =================================== FAILURES ===================================
> _______________ TestCL.test_invalid_kernel_names_cause_failures ________________
>
> self = <test_wrapper.TestCL instance at 0x2622f80>
>
> @pytools.test.mark_test.opencl
> def test_invalid_kernel_names_cause_failures(self):
> for platform in cl.get_platforms():
> for device in platform.get_devices():
> ctx = cl.Context([device])
> prg = cl.Program(ctx, """
> __kernel void sum(__global float *a)
> { a[get_global_id(0)] *= 2; }
> """).build()
>
> try:
> > prg.sam
>
> test_wrapper.py:166:
> _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
>
> self = <pyopencl.Program object at 0x2623950>, attr = 'sam'
>
> def __getattr__(self, attr):
> try:
> > knl = Kernel(self._get_prg(), attr)
> E RuntimeError: No kernel found for given name
Known. This is a bug in the Intel implementation. Funnily, in this case
a C++ exception comes flying past the (supposedly C-only) OpenCL interface.
> ==================================== ERRORS ====================================
> ________________________ ERROR collecting test_array.py ________________________
> test_array.py:665: in <module>
> > register_dtype(mmc_dtype, "minmax_collector")
> ../../../../.local/lib/python2.7/site-packages/pyopencl-2011.1.2-py2.7-linux-x86_64.egg/pyopencl/compyte/dtypes.py:50: in register_dtype
> > % (dtype, DTYPE_TO_NAME[dtype], ", ".join(c_names)))
> E RuntimeError: dtype '[('cur_min', '<f4'), ('cur_max', '<f4')]' already registered (as 'minmax_collector', new names 'minmax_collector')
> =========================== 1 error in 0.01 seconds
> ============================
This should be fixed in recent git.
Andreas
7 years, 1 month
