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
Ghosting with bilinear Interpolation of OpenCV-Images
by Andreas Tharang
Hello
I`ve made a little testscript to test the bilinear interpolation of
image2d_t-objects.
The script:
import pyopencl as cl
import numpy as np
import cv2 # OpenCV 2.3.1
Img = cv2.imread("Test.jpg") # read Image with width = 709px and height
= 472px
Img = cv2.cvtColor(Img, cv2.COLOR_BGR2GRAY) # convert to grayscale
print Img.shape # prints: (472L, 709L)
OutImg = np.empty(shape=Img.shape, dtype=np.uint8) # create Output-Image
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags
dev_Img = cl.Image(ctx,
mf.READ_ONLY | mf.COPY_HOST_PTR,
cl.ImageFormat(cl.channel_order.R,
cl.channel_type.UNSIGNED_INT8),
hostbuf=Img)
dev_OutImg = cl.Image(ctx,
mf.WRITE_ONLY | mf.ALLOC_HOST_PTR,
cl.ImageFormat(cl.channel_order.R,
cl.channel_type.UNSIGNED_INT8),
shape=Img.shape)
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)(0.5f) + (float2)(0.4f)) / convert_float2(dims); // = (x + 0.5 +
0.4)/h ; (y + 0.5 + 0.2)/w
uint4 Pixel = read_imageui(Img, smp, NormCoords);
write_imageui(Out, Coords, Pixel);
}
""").build()
prg.ImageCopy(queue, Img.shape, None, dev_Img, dev_OutImg)
cl.enqueue_read_image(queue, dev_OutImg, (0, 0), OutImg.shape,
OutImg).wait()
cv2.imwrite("Out.jpg", OutImg)
The problem: This is the input-image:
http://s7.directupload.net/file/d/2692/hkbb8pn2_jpg.htm
And when i add a small offset to the image-coords ( here 0.4f) the
output is something like this:
http://s1.directupload.net/file/d/2692/xuqzmtp4_jpg.htm
When i don't add this little offset or use CLK_FILTER_NEAREST then the
result is o.k.
The effect is proportional to the ratio of the image height and width
and the absolute value of the offset.
Also only the offset in image-width direction has a influence ( if i add
(float2)(0.4f,0.0f) the result is nearly o.k. - 4 diagonal pixel-bands
remains)
A quadratic image produces allways the correct output.
The script runs on a NVidia NVS3100M with image-support.
I`ve tried to swap the Image-shapes and Coords in all ways in and
outside the kernel - no good result.
Hope that somebody can guide me back to the right way
Thx Sylvan
8 years, 2 months
Re: [PyOpenCL] invalid work group size, on Mac OSX Lion
by Vincent Favre-Nicolin
Hi Andreas,
----- Mail original -----
"Andreas Kloeckner" <lists(a)informa.tiker.net>:
>
> On Sat, 29 Oct 2011 14:10:39 +0200 (CEST), Vincent Favre-Nicolin
> <vincent.favre-nicolin(a)cea.fr> wrote:
> > return enqueue_nd_range_kernel(queue, self, global_size,
> > local_size,
> > > global_offset, wait_for, g_times_l=g_times_l)
> > E MemoryError: clEnqueueNDRangeKernel failed: mem object
> > allocation failure
>
> The test_random error I saw as well, so that's ok. This latter one is
> also mystifying. What does it think it's allocating here? How much
> memory do you have in your machine? Any chance that you just
> physically ran out of memory?
I doubt it, though the graphic card is not too powerful (this part of the test failed on the "ATI Radeon HD 6630M"). The computer has 4GB of memory (and freshly rebooted), and the card has (according to wikipedia) 1GB of mem.
> Ah, I just had an idea: Can you print local_size right before this
> failure? If it's bigger than 32K, that might be the issue--although
> I'm
> not sure that'd be the case...
global_size,local_size: (2304,) (128,)
Not so big, so... But I've seen a few problems on the ATI card - sometimes I get into a state where opencl initialization becomes impossible, and I have to reboot. Not sure if a soft reset is possible. But that's a different issue.
Vincent
8 years, 3 months
Re: [PyOpenCL] invalid work group size, on Mac OSX Lion
by Andreas Kloeckner
Hi Vincent,
On Sat, 29 Oct 2011 14:10:39 +0200 (CEST), Vincent Favre-Nicolin <vincent.favre-nicolin(a)cea.fr> wrote:
> return enqueue_nd_range_kernel(queue, self, global_size, local_size,
> > global_offset, wait_for, g_times_l=g_times_l)
> E MemoryError: clEnqueueNDRangeKernel failed: mem object allocation
> failure
The test_random error I saw as well, so that's ok. This latter one is
also mystifying. What does it think it's allocating here? How much
memory do you have in your machine? Any chance that you just physically
ran out of memory?
Ah, I just had an idea: Can you print local_size right before this
failure? If it's bigger than 32K, that might be the issue--although I'm
not sure that'd be the case...
Andreas
8 years, 3 months
Re: [PyOpenCL] invalid work group size, on Mac OSX Lion
by Vincent Favre-Nicolin
Hi,
I tried this (current git) on my new Mac mini (running OSX Lion and XCode 4.2), and now test_array has 2 failures (one on the i5, one on the ATI 6630M):
test_array.py ........................F............................F......
cl_math and wrapper passed (one skipped for cl_math).
Below: dump_properties, and result of test_array.
I'm trying my own code and getting issues with "invalid workgroup size" on the i5 (no problem on the ATI card), but it may be my mistake - I still haven't understood everything about work sizes when moving from cuda to opencl...
In the examples, transpose.py fails (in the i5) when starting the "naive" test with an invalid work group size, maybe because of the fixed block size=16 ?
Vincent
####################################################################
Mac-mini-de-Vincent:test vincent$ python ../examples/dump-properties.py
===========================================================================
<pyopencl.Platform 'Apple' at 0x7fff0000>
===========================================================================
EXTENSIONS: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event
NAME: Apple
PROFILE: FULL_PROFILE
VENDOR: Apple
VERSION: OpenCL 1.1 (Jul 25 2011 15:56:07)
Oct 29 13:46:55 Mac-mini-de-Vincent.local Python[144] <Error>: kCGErrorFailure: Set a breakpoint @ CGErrorBreakpoint() to catch errors as they are logged.
---------------------------------------------------------------------------
<pyopencl.Device 'Intel(R) Core(TM) i5-2520M CPU @ 2.50GHz' on 'Apple' at 0xffffffff>
---------------------------------------------------------------------------
ADDRESS_BITS: 64
AVAILABLE: 1
COMPILER_AVAILABLE: 1
DOUBLE_FP_CONFIG: 63
DRIVER_VERSION: 1.1
ENDIAN_LITTLE: 1
ERROR_CORRECTION_SUPPORT: 0
EXECUTION_CAPABILITIES: 3
EXTENSIONS: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_3d_image_writes cl_APPLE_fp64_basic_ops cl_APPLE_fixed_alpha_channel_orders cl_APPLE_biased_fixed_point_image_formats
GLOBAL_MEM_CACHELINE_SIZE: 3145728
GLOBAL_MEM_CACHE_SIZE: 64
GLOBAL_MEM_CACHE_TYPE: 2
GLOBAL_MEM_SIZE: 4294967296
IMAGE2D_MAX_HEIGHT: 8192
IMAGE2D_MAX_WIDTH: 8192
IMAGE3D_MAX_DEPTH: 2048
IMAGE3D_MAX_HEIGHT: 2048
IMAGE3D_MAX_WIDTH: 2048
IMAGE_SUPPORT: 1
LOCAL_MEM_SIZE: 32768
LOCAL_MEM_TYPE: 2
MAX_CLOCK_FREQUENCY: 2500
MAX_COMPUTE_UNITS: 4
MAX_CONSTANT_ARGS: 8
MAX_CONSTANT_BUFFER_SIZE: 65536
MAX_MEM_ALLOC_SIZE: 1073741824
MAX_PARAMETER_SIZE: 4096
MAX_READ_IMAGE_ARGS: 128
MAX_SAMPLERS: 16
MAX_WORK_GROUP_SIZE: 1024
MAX_WORK_ITEM_DIMENSIONS: 3
MAX_WORK_ITEM_SIZES: [1024, 1, 1]
MAX_WRITE_IMAGE_ARGS: 8
MEM_BASE_ADDR_ALIGN: 1024
MIN_DATA_TYPE_ALIGN_SIZE: 128
NAME: Intel(R) Core(TM) i5-2520M CPU @ 2.50GHz
PLATFORM: <pyopencl.Platform 'Apple' at 0x7fff0000>
PREFERRED_VECTOR_WIDTH_CHAR: 16
PREFERRED_VECTOR_WIDTH_DOUBLE: 2
PREFERRED_VECTOR_WIDTH_FLOAT: 4
PREFERRED_VECTOR_WIDTH_INT: 4
PREFERRED_VECTOR_WIDTH_LONG: 2
PREFERRED_VECTOR_WIDTH_SHORT: 8
PROFILE: FULL_PROFILE
PROFILING_TIMER_RESOLUTION: 1
QUEUE_PROPERTIES: 2
SINGLE_FP_CONFIG: 63
TYPE: 2
VENDOR: Intel
VENDOR_ID: 4294967295
VERSION: OpenCL 1.1
IMAGE2D READ_ONLY FORMATS: RGBA-F, <unknown channel order 0x10bc>-F, <unknown channel order 0x10bb>-F, <unknown channel order 0x10ba>-F, INTENSITY-HALF_F, RGBA-HALF_F, <unknown channel order 0x10bc>-HALF_F, RGBA-<unknown channel data type 268435464>, RGBA-S16, RGBA-S32, RGBA-S8, RGBA-SN16, RGBA-SN8, RGBA-U16, RGBA-U32, RGBA-U8, RGBA-UN16, <unknown channel order 0x10bc>-UN16, <unknown channel order 0x10000006>-UN8, A-UN8, <unknown channel order 0x10b7>-UN8, <unknown channel order 0x10000007>-UN8, BGRA-UN8, INTENSITY-UN8, RGBA-UN8, <unknown channel order 0x10bc>-UN8, <unknown channel order 0x10bb>-UN8
IMAGE3D READ_ONLY FORMATS: RGBA-F, <unknown channel order 0x10bc>-F, <unknown channel order 0x10bb>-F, <unknown channel order 0x10ba>-F, INTENSITY-HALF_F, RGBA-HALF_F, <unknown channel order 0x10bc>-HALF_F, RGBA-<unknown channel data type 268435464>, RGBA-S16, RGBA-S32, RGBA-S8, RGBA-SN16, RGBA-SN8, RGBA-U16, RGBA-U32, RGBA-U8, RGBA-UN16, <unknown channel order 0x10bc>-UN16, <unknown channel order 0x10000006>-UN8, A-UN8, <unknown channel order 0x10b7>-UN8, <unknown channel order 0x10000007>-UN8, BGRA-UN8, INTENSITY-UN8, RGBA-UN8, <unknown channel order 0x10bc>-UN8, <unknown channel order 0x10bb>-UN8
---------------------------------------------------------------------------
<pyopencl.Device 'ATI Radeon HD 6630M' on 'Apple' at 0x1021b00>
---------------------------------------------------------------------------
ADDRESS_BITS: 32
AVAILABLE: 1
COMPILER_AVAILABLE: 1
DOUBLE_FP_CONFIG: 0
DRIVER_VERSION: 1.0
ENDIAN_LITTLE: 1
ERROR_CORRECTION_SUPPORT: 0
EXECUTION_CAPABILITIES: 1
EXTENSIONS: cl_APPLE_SetMemObjectDestructor cl_APPLE_ContextLoggingFunctions cl_APPLE_clut cl_APPLE_query_kernel_names cl_APPLE_gl_sharing cl_khr_gl_event cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_3d_image_writes
GLOBAL_MEM_CACHELINE_SIZE: 0
GLOBAL_MEM_CACHE_SIZE: 0
GLOBAL_MEM_CACHE_TYPE: 0
GLOBAL_MEM_SIZE: 134217728
IMAGE2D_MAX_HEIGHT: 8192
IMAGE2D_MAX_WIDTH: 8192
IMAGE3D_MAX_DEPTH: 2048
IMAGE3D_MAX_HEIGHT: 2048
IMAGE3D_MAX_WIDTH: 2048
IMAGE_SUPPORT: 1
LOCAL_MEM_SIZE: 32768
LOCAL_MEM_TYPE: 1
MAX_CLOCK_FREQUENCY: 485
MAX_COMPUTE_UNITS: 6
MAX_CONSTANT_ARGS: 8
MAX_CONSTANT_BUFFER_SIZE: 65536
MAX_MEM_ALLOC_SIZE: 134217728
MAX_PARAMETER_SIZE: 1024
MAX_READ_IMAGE_ARGS: 128
MAX_SAMPLERS: 16
MAX_WORK_GROUP_SIZE: 1024
MAX_WORK_ITEM_DIMENSIONS: 3
MAX_WORK_ITEM_SIZES: [1024, 1024, 1024]
MAX_WRITE_IMAGE_ARGS: 8
MEM_BASE_ADDR_ALIGN: 32768
MIN_DATA_TYPE_ALIGN_SIZE: 128
NAME: ATI Radeon HD 6630M
PLATFORM: <pyopencl.Platform 'Apple' at 0x7fff0000>
PREFERRED_VECTOR_WIDTH_CHAR: 16
PREFERRED_VECTOR_WIDTH_DOUBLE: 0
PREFERRED_VECTOR_WIDTH_FLOAT: 4
PREFERRED_VECTOR_WIDTH_INT: 4
PREFERRED_VECTOR_WIDTH_LONG: 2
PREFERRED_VECTOR_WIDTH_SHORT: 8
PROFILE: FULL_PROFILE
PROFILING_TIMER_RESOLUTION: 37
QUEUE_PROPERTIES: 2
SINGLE_FP_CONFIG: 30
TYPE: 4
VENDOR: AMD
VENDOR_ID: 16915200
VERSION: OpenCL 1.1
IMAGE2D READ_ONLY FORMATS: RGBA-SN8, RGBA-SN16, RGBA-UN8, RGBA-UN16, RGBA-S8, RGBA-S16, RGBA-S32, RGBA-U8, RGBA-U16, RGBA-U32, RGBA-HALF_F, RGBA-F, BGRA-UN8, <unknown channel order 0x10b7>-UN8, R-SN8, R-UN8, R-UN16, R-S8, R-S16, R-S32, R-U8, R-U16, R-U32, R-HALF_F, R-F, A-SN8, A-UN8, A-UN16, A-S8, A-S16, A-S32, A-U8, A-U16, A-U32, A-HALF_F, A-F, RG-UN8, RG-UN16, RG-S8, RG-S16, RG-S32, RG-U8, RG-U16, RG-U32, RG-HALF_F, RG-F, RA-UN8, RA-UN16, RA-S8, RA-S16, RA-S32, RA-U8, RA-U16, RA-U32, RA-HALF_F, RA-F, INTENSITY-UN8, INTENSITY-UN16, INTENSITY-HALF_F, INTENSITY-F
IMAGE3D READ_ONLY FORMATS: RGBA-SN8, RGBA-SN16, RGBA-UN8, RGBA-UN16, RGBA-S8, RGBA-S16, RGBA-S32, RGBA-U8, RGBA-U16, RGBA-U32, RGBA-HALF_F, RGBA-F, BGRA-UN8, <unknown channel order 0x10b7>-UN8, R-SN8, R-UN8, R-UN16, R-S8, R-S16, R-S32, R-U8, R-U16, R-U32, R-HALF_F, R-F, A-SN8, A-UN8, A-UN16, A-S8, A-S16, A-S32, A-U8, A-U16, A-U32, A-HALF_F, A-F, RG-UN8, RG-UN16, RG-S8, RG-S16, RG-S32, RG-U8, RG-U16, RG-U32, RG-HALF_F, RG-F, RA-UN8, RA-UN16, RA-S8, RA-S16, RA-S32, RA-U8, RA-U16, RA-U32, RA-HALF_F, RA-F, INTENSITY-UN8, INTENSITY-UN16, INTENSITY-HALF_F, INTENSITY-F
####################################################################
Mac-mini-de-Vincent:test vincent$ python test_array.py
================================================================================ test session starts ================================================================================
platform darwin -- Python 2.7.2 -- pytest-2.1.3
collected 60 items
test_array.py ........................F............................F......
===================================================================================== FAILURES ======================================================================================
_______________________ test_random[ctx_factory=<context factory for <pyopencl.Device 'Intel(R) Core(TM) i5-2520M CPU @ 2.50GHz' on 'Apple' at 0xffffffff>>] ________________________
ctx_factory = <pyopencl.tools.ContextFactory instance at 0x1059a4680>
@pytools.test.mark_test.opencl
def test_random(ctx_factory):
context = ctx_factory()
queue = cl.CommandQueue(context)
from pyopencl.clrandom import RanluxGenerator
if has_double_support(context.devices[0]):
dtypes = [np.float32, np.float64]
else:
dtypes = [np.float32]
gen = RanluxGenerator(queue, 5120)
for ary_size in [300, 301, 302, 303, 10007]:
for dtype in dtypes:
ran = cl_array.zeros(queue, ary_size, dtype)
gen.fill_uniform(ran)
assert (0 < ran.get()).all()
assert (ran.get() < 1).all()
gen.synchronize(queue)
ran = cl_array.zeros(queue, ary_size, dtype)
gen.fill_uniform(ran, a=4, b=7)
> assert (4 < ran.get()).all()
E assert <built-in method all of numpy.ndarray object at 0x7f82e4d006d0>()
E + where <built-in method all of numpy.ndarray object at 0x7f82e4d006d0> = 4 < array([ 6.19607544, 4.17939472, 4.86412287, 6.28594398, 4. ,\n ...4. , 4. , 4. , 4. , 4. ], dtype=float32).all
E + where array([ 6.19607544, 4.17939472, 4.86412287, 6.28594398, 4. ,\n ...4. , 4. , 4. , 4. , 4. ], dtype=float32) = <bound method Array.get of array([ 6.19607544, 4.17939472, 4.86412287, 6.28.... , 4. , 4. , 4. , 4. ], dtype=float32)>()
E + where <bound method Array.get of array([ 6.19607544, 4.17939472, 4.86412287, 6.28.... , 4. , 4. , 4. , 4. ], dtype=float32)> = array([ 6.19607544, 4.17939472, 4.86412287, 6.28594398, 4. ,\n ...4. , 4. , 4. , 4. , 4. ], dtype=float32).get
test_array.py:249: AssertionError
---------------------------------------------------------------------------------- Captured stderr ----------------------------------------------------------------------------------
/opt/local/Library/Frameworks/Python.framework/Versions/2.7/lib/python2.7/site-packages/pyopencl-2011.1.2-py2.7-macosx-10.7-x86_64.egg/pyopencl/__init__.py:36: CompilerWarning: Non-empty compiler output encountered. Set the environment variable PYOPENCL_COMPILER_OUTPUT=1 to see more.
"to see more.", CompilerWarning)
___________________________________ test_scan[ctx_factory=<context factory for <pyopencl.Device 'ATI Radeon HD 6630M' on 'Apple' at 0x1021b00>>] ____________________________________
ctx_factory = <pyopencl.tools.ContextFactory instance at 0x1059ab248>
@pytools.test.mark_test.opencl
def test_scan(ctx_factory):
context = ctx_factory()
queue = cl.CommandQueue(context)
from pyopencl.scan import InclusiveScanKernel, ExclusiveScanKernel
dtype = np.int32
for cls in [InclusiveScanKernel, ExclusiveScanKernel]:
knl = cls(context, dtype, "a+b", "0")
for n in [
10, 2 ** 10 - 5, 2 ** 10,
2 ** 20 - 2 ** 18,
2 ** 20 - 2 ** 18 + 5,
2 ** 10 + 5,
2 ** 20 + 5,
2 ** 20, 2 ** 24]:
host_data = np.random.randint(0, 10, n).astype(dtype)
dev_data = cl_array.to_device(queue, host_data)
> knl(dev_data)
test_array.py:592:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
self = <pyopencl.compyte.scan.InclusiveScanKernel object at 0x1062e4650>, input_ary = array([9, 2, 8, ..., 9, 9, 0], dtype=int32)
output_ary = array([9, 2, 8, ..., 9, 9, 0], dtype=int32), allocator = <pyopencl._cl.CLAllocator object at 0x106362788>, queue = <pyopencl._cl.CommandQueue object at 0x106362680>
def __call__(self, input_ary, output_ary=None, allocator=None,
queue=None):
allocator = allocator or input_ary.allocator
queue = queue or input_ary.queue or output_ary.queue
if output_ary is None:
output_ary = input_ary
if isinstance(output_ary, (str, unicode)) and output_ary == "new":
output_ary = cl_array.empty_like(input_ary, allocator=allocator)
if input_ary.shape != output_ary.shape:
raise ValueError("input and output must have the same shape")
if not input_ary.flags.forc:
raise RuntimeError("ScanKernel cannot "
"deal with non-contiguous arrays")
n, = input_ary.shape
if not n:
return output_ary
unit_size = self.scan_wg_size * self.scan_wg_seq_batches
max_groups = 3*max(dev.max_compute_units for dev in self.devices)
from pytools import uniform_interval_splitting
interval_size, num_groups = uniform_interval_splitting(
n, unit_size, max_groups);
block_results = allocator(self.dtype.itemsize*num_groups)
dummy_results = allocator(self.dtype.itemsize)
# first level scan of interval (one interval per block)
self.scan_intervals_knl(
queue, (num_groups*self.scan_wg_size,), (self.scan_wg_size,),
input_ary.data,
n, interval_size,
output_ary.data,
> block_results)
/opt/local/Library/Frameworks/Python.framework/Versions/2.7/lib/python2.7/site-packages/pyopencl-2011.1.2-py2.7-macosx-10.7-x86_64.egg/pyopencl/compyte/scan.py:481:
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _
self = <pyopencl._cl.Kernel object at 0x105974d08>, queue = <pyopencl._cl.CommandQueue object at 0x106362680>, global_size = (2304,)
def kernel_call(self, queue, global_size, *args, **kwargs):
global_offset = kwargs.pop("global_offset", None)
had_local_size = "local_size" in kwargs
local_size = kwargs.pop("local_size", None)
g_times_l = kwargs.pop("g_times_l", False)
wait_for = kwargs.pop("wait_for", None)
if kwargs:
raise TypeError(
"Kernel.__call__ recived unexpected keyword arguments: %s"
% ", ".join(kwargs.keys()))
if had_local_size:
from warnings import warn
warn("The local_size keyword argument is deprecated and will be "
"removed in pyopencl 2012.x. Pass the local "
"size as the third positional argument instead.",
DeprecationWarning, stacklevel=2)
if isinstance(args[0], (type(None), tuple)) and not had_local_size:
local_size = args[0]
args = args[1:]
elif not had_local_size:
from warnings import warn
warn("PyOpenCL Warning: There was an API change "
"in Kernel.__call__() in pyopencl 0.92. "
"local_size was moved from keyword argument to third "
"positional argument in pyopencl 0.92. "
"You didn't pass local_size, but you still need to insert "
"'None' as a third argument. "
"Your present usage is deprecated and will stop "
"working in pyopencl 2012.x.",
DeprecationWarning, stacklevel=2)
self.set_args(*args)
return enqueue_nd_range_kernel(queue, self, global_size, local_size,
> global_offset, wait_for, g_times_l=g_times_l)
E MemoryError: clEnqueueNDRangeKernel failed: mem object allocation failure
/opt/local/Library/Frameworks/Python.framework/Versions/2.7/lib/python2.7/site-packages/pyopencl-2011.1.2-py2.7-macosx-10.7-x86_64.egg/pyopencl/__init__.py:247: MemoryError
======================================================================= 2 failed, 58 passed in 15.14 seconds ========================================================================
8 years, 3 months
Re: [PyOpenCL] invalid work group size
by Andreas Kloeckner
Hi Yosuke,
On Sun, 25 Sep 2011 21:22:04 +0900, onoue(a)likr-lab.com wrote:
> I found a Apple OpenCL bug on Mac OS X Lion. If device type is CPU,
> device.max_work_item_sizes[0] returns 1024. But, invalid work group
> size error is occurred when local_size[0] is greater than 128. And, I
> don't know why, ElementwiseKernel works only when local_size is None
> or (1, 1, 1). I have already reported this error to Apple.
>
> Incidentally, I got following error when I execute a script in Apple
> OpenCL with PyOpenCL source code getting from git repository.
>
> Traceback (most recent call last):
> File "montecarlo.py", line 27, in <module>
> kernel = create_kernel(context, D)
> File "montecarlo.py", line 17, in create_kernel
> arguments=arguments)
> File "/Users/likr/git/pyopencl/pyopencl/reduction.py", line 268, in __init__
> name=name+"_stage1", options=options, preamble=preamble)
> File "/Users/likr/git/pyopencl/pyopencl/reduction.py", line 234, in
> get_reduction_kernel
> name, preamble, device, max_group_size)
> File "/Users/likr/git/pyopencl/pyopencl/reduction.py", line 185, in
> get_reduction_source
> no_sync_size = min(get_dev_no_sync_size(dev) for dev in devices)
> File "/Users/likr/git/pyopencl/pyopencl/reduction.py", line 185, in <genexpr>
> no_sync_size = min(get_dev_no_sync_size(dev) for dev in devices)
> File "/Users/likr/git/pyopencl/pyopencl/reduction.py", line 175, in
> get_dev_no_sync_size
> result = get_simd_group_size(device)
> File "/Users/likr/git/pyopencl/pyopencl/characterize.py", line 271,
> in get_simd_group_size
> if dtype.itemsize == 1:
> NameError: global name 'dtype' is not defined
This should not be an issue any more with the current git version of
PyOpenCL, soon to be PyOpenCL 2011.2.
Thanks for the report, and apologies for the long delay.
Andreas
8 years, 3 months
clearing kernel cache
by andrew cooke
i am having trouble working out how to clear the kernel cache. the docs
suggest clear_context_caches() from pyopencl.tools, but it doesn't seem to be
there:
Python 3.2 (r32:88445, Jun 9 2011, 09:31:13)
[GCC 4.5.1 20101208 [gcc-4_5-branch revision 167585]] on linux2
Type "help", "copyright", "credits" or "license" for more information.
>>> from pyopencl.tools import clear_context_caches
Traceback (most recent call last):
File "<stdin>", line 1, in <module>
ImportError: cannot import name clear_context_caches
>>> from pyopencl import VERSION
>>> VERSION
(2011, 1, 2)
what should i be doing?
thanks,
andrew
8 years, 3 months
Aparapi like features
by Ziyad Saeed
AMD released Aparapi project that allows Java developers to use GPU and does it in a wonderful way.
First you don't right C code at all. It converts the java code at runtime to be able to run on the GPU. PyOpenCL does that as well via meta programming. check
Second it is smart enough to adapt itself from a single code whether to execute on a OpenCL device or not. You never write two separate codes. If an OpenCL device is detected it uses it, if not then uses java thread pool. Essentially taking all the headache away. It automatically covers all your bases with a single code base.
Can PyOpenCL adapt itself if OpenCL device isn't detected.
8 years, 3 months
Re: [PyOpenCL] Problem starting
by Andreas Kloeckner
Hi Pedro,
On Mon, 3 Oct 2011 23:10:46 -0700, Pedro Marcal <pedrovmarcal(a)gmail.com> wrote:
> I installed the binary from the Irvine site, but cannot get the included
> example to execute.
> It looks like I need some path or a .dll is missing.
> Here is the diagnostic for the starting few lines
>
> import pyopencl as cl
> import numpy
> import numpy.linalg as la
>
> resulting in the following error report.
>
> Traceback (most recent call last):
> File "C:\Python26\OpenCl_ex.py", line 1, in <module>
> import pyopencl as cl
> File "C:\Python26\lib\site-packages\pyopencl\__init__.py", line 4, in
> <module>
> import pyopencl._cl as _cl
> ImportError: DLL load failed: The specified module could not be found.
Check the directory C:\Python26\lib\site-packages\pyopencl\ and look for
a file named _cl.pyd or _cl.dll. Then use the tool 'Dependency Walker'
to check whether it can find all the DLLs it needs.
HTH,
Andreas
8 years, 3 months
