diff --git a/pyopencl/cffi_cl.py b/pyopencl/cffi_cl.py
index 356c64b020d4cdc3feecaecf1696240f741e5a94..c2e87ca1035bfa1e6363ea43a2e6c5205c218a72 100644
--- a/pyopencl/cffi_cl.py
+++ b/pyopencl/cffi_cl.py
@@ -1862,7 +1862,7 @@ class Kernel(_Common):
from pyopencl.characterize import has_struct_arg_count_bug
count_bug_per_dev = [
- has_struct_arg_count_bug(dev)
+ has_struct_arg_count_bug(dev, self.context)
for dev in self.context.devices]
from pytools import single_valued
diff --git a/pyopencl/characterize/__init__.py b/pyopencl/characterize/__init__.py
index b6ea1c53f4d812dd79389e17a11051d88a4be173..d03051897b8e35a1a114fa5e8f4cebd145bb1589 100644
--- a/pyopencl/characterize/__init__.py
+++ b/pyopencl/characterize/__init__.py
@@ -1,8 +1,4 @@
-from __future__ import division
-from __future__ import absolute_import
-import six
-from six.moves import range
-from six.moves import zip
+from __future__ import division, absolute_import
__copyright__ = "Copyright (C) 2009 Andreas Kloeckner"
@@ -28,6 +24,8 @@ THE SOFTWARE.
import pyopencl as cl
from pytools import memoize
+import six
+from six.moves import range, zip
class CLCharacterizationWarning(UserWarning):
@@ -322,14 +320,70 @@ def get_simd_group_size(dev, type_size):
return None
-def has_struct_arg_count_bug(dev):
+def get_pocl_version(platform, fallback_value=None):
+ if platform.name != "Portable Computing Language":
+ return None
+
+ import re
+ ver_match = re.match(
+ r"^OpenCL [0-9.]+ pocl ([0-9]+)\.([0-9]+)", platform.version)
+ if ver_match is None:
+ msg = ("pocl version number did not have expected format: '%s'"
+ % platform.version)
+ if fallback_value is not None:
+ from warnings import warn
+ warn(msg)
+ return fallback_value
+ else:
+ raise ValueError(msg)
+ else:
+ return (int(ver_match.group(1)), int(ver_match.group(2)))
+
+
+_CHECK_FOR_POCL_ARG_COUNT_BUG_CACHE = {}
+
+
+def _check_for_pocl_arg_count_bug(dev, ctx=None):
+ try:
+ return _CHECK_FOR_POCL_ARG_COUNT_BUG_CACHE[dev]
+ except KeyError:
+ pass
+
+ if ctx is None:
+ build_ctx = cl.Context([dev])
+ else:
+ build_ctx = ctx
+
+ prg = cl.Program(build_ctx, """
+ struct two_things
+ {
+ long a;
+ long b;
+ };
+
+ __kernel void test_knl(struct two_things x)
+ {
+ }
+ """).build()
+
+ result = prg.test_knl.num_args == 2
+ _CHECK_FOR_POCL_ARG_COUNT_BUG_CACHE[dev] = result
+
+ return result
+
+
+def has_struct_arg_count_bug(dev, ctx=None):
"""Checks whether the device is expected to have the
`argument counting bug <https://github.com/pocl/pocl/issues/197>`_.
"""
if dev.platform.name == "Apple" and dev.type & cl.device_type.CPU:
return "apple"
- if (dev.platform.name == "Portable Computing Language"
- and dev.address_bits == 64):
- return "pocl"
+ if dev.platform.name == "Portable Computing Language":
+ pocl_version = get_pocl_version(dev.platform, fallback_value=(0.14))
+ if pocl_version <= (0, 13):
+ return "pocl"
+ elif pocl_version <= (0, 14) and _check_for_pocl_arg_count_bug(dev, ctx):
+ return "pocl"
+
return False
diff --git a/test/test_algorithm.py b/test/test_algorithm.py
index 374381ede72ad66951d5146613f4f7a00cc9311a..23ebe4ef184c536657578687edfb2300970acd52 100644
--- a/test/test_algorithm.py
+++ b/test/test_algorithm.py
@@ -386,7 +386,8 @@ def test_dot(ctx_factory):
vdot_ab_gpu = cl_array.vdot(a_gpu, b_gpu).get()
- assert abs(vdot_ab_gpu - vdot_ab) / abs(vdot_ab) < 1e-4
+ rel_err = abs(vdot_ab_gpu - vdot_ab) / abs(vdot_ab)
+ assert rel_err < 1e-4, rel_err
@memoize
diff --git a/test/test_enqueue_copy.py b/test/test_enqueue_copy.py
index 14f0bc7df3191531099b6077754c0c77870ffda8..564e833a4bc167644bd057c5fa2117d17bda2cda 100644
--- a/test/test_enqueue_copy.py
+++ b/test/test_enqueue_copy.py
@@ -29,6 +29,7 @@ import pytest
from pyopencl.tools import ( # noqa
pytest_generate_tests_for_pyopencl as pytest_generate_tests)
+from pyopencl.characterize import get_pocl_version
def generate_slice(start, shape):
@@ -42,7 +43,9 @@ def test_enqueue_copy_rect_2d(ctx_factory, honor_skip=True):
ctx = ctx_factory()
queue = cl.CommandQueue(ctx)
- if honor_skip and ctx.devices[0].platform.name == "Portable Computing Language":
+ if (honor_skip
+ and ctx.devices[0].platform.name == "Portable Computing Language"
+ and get_pocl_version(ctx.devices[0].platform) <= (0, 13)):
# https://github.com/pocl/pocl/issues/353
pytest.skip("POCL's rectangular copies crash")
@@ -127,7 +130,9 @@ def test_enqueue_copy_rect_3d(ctx_factory, honor_skip=True):
ctx = ctx_factory()
queue = cl.CommandQueue(ctx)
- if honor_skip and ctx.devices[0].platform.name == "Portable Computing Language":
+ if (honor_skip
+ and ctx.devices[0].platform.name == "Portable Computing Language"
+ and get_pocl_version(ctx.devices[0].platform) <= (0, 13)):
# https://github.com/pocl/pocl/issues/353
pytest.skip("POCL's rectangular copies crash")
diff --git a/test/test_wrapper.py b/test/test_wrapper.py
index b1a23c2a00eb2289fd98dd5a228953109bf3ffef..a281e3105c5d17941821de6c103609a298d99a86 100644
--- a/test/test_wrapper.py
+++ b/test/test_wrapper.py
@@ -33,6 +33,7 @@ import pyopencl.array as cl_array
import pyopencl.clrandom
from pyopencl.tools import ( # noqa
pytest_generate_tests_for_pyopencl as pytest_generate_tests)
+from pyopencl.characterize import get_pocl_version
# Are CL implementations crashy? You be the judge. :)
try:
@@ -43,10 +44,11 @@ else:
faulthandler.enable()
-def _skip_if_pocl(plat, msg='unsupported by pocl'):
+def _skip_if_pocl(plat, up_to_version, msg='unsupported by pocl'):
if plat.vendor == "The pocl project":
- import pytest
- pytest.skip(msg)
+ if up_to_version is None or get_pocl_version(plat) <= up_to_version:
+ import pytest
+ pytest.skip(msg)
def test_get_info(ctx_factory):
@@ -363,7 +365,7 @@ def test_image_2d(ctx_factory):
if "Intel" in device.vendor and "31360.31426" in device.version:
from pytest import skip
skip("images crashy on %s" % device)
- _skip_if_pocl(device.platform, 'pocl does not support CL_ADDRESS_CLAMP')
+ _skip_if_pocl(device.platform, None, 'pocl does not support CL_ADDRESS_CLAMP')
prg = cl.Program(context, """
__kernel void copy_image(
@@ -435,7 +437,7 @@ def test_image_3d(ctx_factory):
if device.platform.vendor == "Intel(R) Corporation":
from pytest import skip
skip("images crashy on %s" % device)
- _skip_if_pocl(device.platform, 'pocl does not support CL_ADDRESS_CLAMP')
+ _skip_if_pocl(device.platform, None, 'pocl does not support CL_ADDRESS_CLAMP')
prg = cl.Program(context, """
__kernel void copy_image_plane(
@@ -626,7 +628,8 @@ def test_can_build_binary(ctx_factory):
def test_enqueue_barrier_marker(ctx_factory):
ctx = ctx_factory()
- _skip_if_pocl(ctx.devices[0].platform, 'pocl crashes on enqueue_barrier')
+ # Still relevant on pocl 0.14.
+ _skip_if_pocl(ctx.devices[0].platform, None, 'pocl crashes on enqueue_barrier')
queue = cl.CommandQueue(ctx)
cl.enqueue_barrier(queue)
evt1 = cl.enqueue_marker(queue)
@@ -647,7 +650,7 @@ def test_unload_compiler(platform):
cl.get_cl_header_version() < (1, 2)):
from pytest import skip
skip("clUnloadPlatformCompiler is only available in OpenCL 1.2")
- _skip_if_pocl(platform, 'pocl does not support unloading compiler')
+ _skip_if_pocl(platform, (0, 13), 'pocl does not support unloading compiler')
if platform.vendor == "Intel(R) Corporation":
from pytest import skip
skip("Intel proprietary driver does not support unloading compiler")
@@ -954,7 +957,8 @@ def test_coarse_grain_svm(ctx_factory):
# https://bitbucket.org/pypy/numpy/issues/52
assert isinstance(svm_ary.mem.base, cl.SVMAllocation)
- if dev.platform.name != "Portable Computing Language":
+ if (dev.platform.name != "Portable Computing Language"
+ or get_pocl_version(dev.platform) >= (0, 14)):
# pocl 0.13 has a bug misinterpreting the size parameter
cl.enqueue_svm_memfill(queue, svm_ary, np.zeros((), svm_ary.mem.dtype))
@@ -980,7 +984,7 @@ def test_coarse_grain_svm(ctx_factory):
if ctx.devices[0].platform.name != "Portable Computing Language":
# "Blocking memcpy is unimplemented (clEnqueueSVMMemcpy.c:61)"
- # in pocl 0.13.
+ # in pocl 0.13 and 0.14-pre.
cl.enqueue_copy(queue, new_ary, svm_ary)
assert np.array_equal(orig_ary*2, new_ary)