From 51b090211b998e0410d471bbd77df8038778601b Mon Sep 17 00:00:00 2001
From: Andreas Kloeckner <inform@tiker.net>
Date: Mon, 20 Feb 2017 23:36:01 -0600
Subject: [PATCH] Rework pocl arg counting workaround to actually test for
presence of bug
---
pyopencl/cffi_cl.py | 2 +-
pyopencl/characterize/__init__.py | 62 ++++++++++++++++++++++++++-----
2 files changed, 54 insertions(+), 10 deletions(-)
diff --git a/pyopencl/cffi_cl.py b/pyopencl/cffi_cl.py
index 45a0ab01..c0dd1b1c 100644
--- a/pyopencl/cffi_cl.py
+++ b/pyopencl/cffi_cl.py
@@ -1850,7 +1850,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 b6ea1c53..1c3d3920 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,60 @@ def get_simd_group_size(dev, type_size):
return None
-def has_struct_arg_count_bug(dev):
+_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":
+ import re
+ ver_match = re.match(
+ r"^OpenCL [0-9.]+ pocl ([0-9]+)\.([0-9]+)", dev.platform.version)
+ if ver_match is None:
+ from warnings import warn
+ warn("pocl version number did not have expected format: '%s'"
+ % dev.platform.version)
+ pocl_version = (0, 14)
+ else:
+ pocl_version = (int(ver_match.group(1)), int(ver_match.group(2)))
+
+ 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
--
GitLab