diff --git a/pyopencl/clmath.py b/pyopencl/clmath.py
new file mode 100644
index 0000000000000000000000000000000000000000..9500486c01003a6b420b2cdae8facf21bac5ca3d
--- /dev/null
+++ b/pyopencl/clmath.py
@@ -0,0 +1,151 @@
+import pyopencl.array as cl_array
+import pyopencl.elementwise as elementwise
+
+def _make_unary_array_func(name):
+ def f(array, queue=None):
+ result = array._new_like_me()
+
+ knl = elementwise.get_unary_func_kernel(array.context, name, array.dtype)
+ knl(queue or array.queue, array._global_size, array._local_size,
+ array.data, result.data, array.mem_size)
+
+ return result
+ return f
+
+# See table 6.8 in the CL spec
+acos = _make_unary_array_func("acos")
+acosh = _make_unary_array_func("acosh")
+acospi = _make_unary_array_func("acospi")
+
+asin = _make_unary_array_func("asin")
+asinh = _make_unary_array_func("asinh")
+asinpi = _make_unary_array_func("asinpi")
+
+atan = _make_unary_array_func("atan")
+# TODO: atan2
+atanh = _make_unary_array_func("atanh")
+atanpi = _make_unary_array_func("atanpi")
+# TODO: atan2pi
+
+cbrt = _make_unary_array_func("cbrt")
+ceil = _make_unary_array_func("ceil")
+# TODO: copysign
+
+cos = _make_unary_array_func("cos")
+cosh = _make_unary_array_func("cosh")
+cospi = _make_unary_array_func("cospi")
+
+erfc = _make_unary_array_func("erfc")
+erf = _make_unary_array_func("erf")
+exp = _make_unary_array_func("exp")
+exp2 = _make_unary_array_func("exp2")
+exp10 = _make_unary_array_func("exp10")
+expm1 = _make_unary_array_func("expm1")
+
+fabs = _make_unary_array_func("fabs")
+# TODO: fdim
+floor = _make_unary_array_func("floor")
+# TODO: fma
+# TODO: fmax
+# TODO: fmin
+
+def fmod(arg, mod, queue=None):
+ """Return the floating point remainder of the division `arg/mod`,
+ for each element in `arg` and `mod`."""
+ result = arg._new_like_me()
+
+ knl = elementwise.get_fmod_kernel(arg.context)
+ knl(queue or arg.queue, arg._global_size, arg._local_size,
+ arg.data, mod.data, result.data, arg.mem_size)
+
+ return result
+
+# TODO: fract
+
+def frexp(arg, queue=None):
+ """Return a tuple `(significands, exponents)` such that
+ `arg == significand * 2**exponent`.
+ """
+ sig = arg._new_like_me()
+ expt = arg._new_like_me()
+
+ knl = elementwise.get_frexp_kernel(arg.context)
+ knl(queue or arg.queue, arg._global_size, arg._local_size,
+ arg.data, sig.data, expt.data, arg.mem_size)
+
+ return sig, expt
+
+# TODO: hypot
+
+ilogb = _make_unary_array_func("ilogb")
+
+def ldexp(significand, exponent, queue=None):
+ """Return a new array of floating point values composed from the
+ entries of `significand` and `exponent`, paired together as
+ `result = significand * 2**exponent`.
+ """
+ result = significand._new_like_me()
+
+ knl = elementwise.get_ldexp_kernel(significand.context)
+ knl(queue or significand.queue,
+ significand._global_size, significand._local_size,
+ significand.data, exponent.data, result.data,
+ significand.mem_size)
+
+ return result
+
+lgamma = _make_unary_array_func("lgamma")
+# TODO: lgamma_r
+
+log = _make_unary_array_func("log")
+log2 = _make_unary_array_func("log2")
+log10 = _make_unary_array_func("log10")
+log1p = _make_unary_array_func("log1p")
+logb = _make_unary_array_func("logb")
+
+# TODO: mad
+# TODO: maxmag
+# TODO: minmag
+
+def modf(arg, queue=None):
+ """Return a tuple `(fracpart, intpart)` of arrays containing the
+ integer and fractional parts of `arg`.
+ """
+
+ intpart = arg._new_like_me()
+ fracpart = arg._new_like_me()
+
+ knl = elementwise.get_modf_kernel(arg.context)
+ knl(queue or arg.queue, arg._global_size, arg._local_size,
+ arg.data, intpart.data, fracpart.data,
+ arg.mem_size)
+
+ return fracpart, intpart
+
+nan = _make_unary_array_func("nan")
+
+# TODO: nextafter
+# TODO: remainder
+# TODO: remquo
+
+rint = _make_unary_array_func("rint")
+# TODO: rootn
+round = _make_unary_array_func("round")
+
+sin = _make_unary_array_func("sin")
+# TODO: sincos
+sinh = _make_unary_array_func("sinh")
+sinpi = _make_unary_array_func("sinpi")
+
+sqrt = _make_unary_array_func("sqrt")
+
+tan = _make_unary_array_func("tan")
+tanh = _make_unary_array_func("tanh")
+tanpi = _make_unary_array_func("tanpi")
+tgamma = _make_unary_array_func("tgamma")
+trunc = _make_unary_array_func("trunc")
+
+# no point wrapping half_ or native_
+
+# TODO: table 6.10, integer functions
+# TODO: table 6.12, clamp et al
diff --git a/pyopencl/elementwise.py b/pyopencl/elementwise.py
index 1a3f25ec60ab757d25e956bf4d071e494a585d40..0fc9aaedf2226fffb03c21745846033c3ee28bca 100644
--- a/pyopencl/elementwise.py
+++ b/pyopencl/elementwise.py
@@ -435,7 +435,7 @@ def get_frexp_kernel(context):
def get_ldexp_kernel(context):
return get_elwise_kernel(context,
"float *sig, float *expt, float *z",
- "z[i] = ldexp(sig[i], int(expt[i]))",
+ "z[i] = ldexp(sig[i], (int) expt[i])",
"ldexp_kernel")
@context_dependent_memoize
diff --git a/pyopencl/tools.py b/pyopencl/tools.py
index aca8c4c91f088aa9a3e73815718322b3b5073ee1..15ade40d93bba3192590a6cff7dd1f00bc0e4a37 100644
--- a/pyopencl/tools.py
+++ b/pyopencl/tools.py
@@ -31,6 +31,7 @@ OTHER DEALINGS IN THE SOFTWARE.
import numpy
from decorator import decorator
+import pyopencl as cl
@@ -57,6 +58,45 @@ def context_dependent_memoize(func, context, *args):
+
+def pytest_generate_tests_for_pyopencl(metafunc):
+ class ContextGetter:
+ def __init__(self, device):
+ self.device = device
+
+ def __call__(self):
+ return cl.Context([device])
+
+ def __str__(self):
+ return "<context getter for %s>" % self.device
+ if ("device" in metafunc.funcargnames
+ or "ctx_getter" in metafunc.funcargnames):
+ arg_dict = {}
+
+ for platform in cl.get_platforms():
+ if "platform" in metafunc.funcargnames:
+ arg_dict["platform"] = platform
+
+ for device in platform.get_devices():
+ if "device" in metafunc.funcargnames:
+ arg_dict["device"] = device
+
+ if "ctx_getter" in metafunc.funcargnames:
+ arg_dict["ctx_getter"] = ContextGetter(device)
+
+ metafunc.addcall(funcargs=arg_dict.copy(),
+ id=", ".join("%s=%s" % (arg, value)
+ for arg, value in arg_dict.iteritems()))
+
+ elif "platform" in metafunc.funcargnames:
+ for platform in cl.get_platforms():
+ metafunc.addcall(
+ funcargs=dict(platform=platform),
+ id=str(platform))
+
+
+
+
# {{{ C code generation helpers -----------------------------------------------
def dtype_to_ctype(dtype):
if dtype is None:
diff --git a/test/test_array.py b/test/test_array.py
index a1f9d2634af01fdc136a6b8fe5efff1ccdb6049c..1fe64d6b7a353d32142399a42eeaff3fc54b4f8a 100644
--- a/test/test_array.py
+++ b/test/test_array.py
@@ -17,6 +17,8 @@ def have_cl():
if have_cl():
import pyopencl.array as cl_array
import pyopencl as cl
+ from pyopencl.tools import pytest_generate_tests_for_pyopencl \
+ as pytest_generate_tests
@@ -546,41 +548,6 @@ def test_astype(ctx_getter):
-def pytest_generate_tests(metafunc):
- if have_cl():
- import pyopencl as cl
- else:
- # will still show "skipped" messages
- return
-
- if ("device" in metafunc.funcargnames
- or "ctx_getter" in metafunc.funcargnames):
- arg_dict = {}
-
- for platform in cl.get_platforms():
- if "platform" in metafunc.funcargnames:
- arg_dict["platform"] = platform
-
- for device in platform.get_devices():
- if "device" in metafunc.funcargnames:
- arg_dict["device"] = device
-
- if "ctx_getter" in metafunc.funcargnames:
- arg_dict["ctx_getter"] = lambda: cl.Context([device])
-
- metafunc.addcall(funcargs=arg_dict.copy(),
- id=", ".join("%s=%s" % (arg, value)
- for arg, value in arg_dict.iteritems()))
-
- elif "platform" in metafunc.funcargnames:
- for platform in cl.get_platforms():
- metafunc.addcall(
- funcargs=dict(platform=platform),
- id=str(platform))
-
-
-
-
if __name__ == "__main__":
# make sure that import failures get reported, instead of skipping the tests.
import pyopencl as cl
diff --git a/test/test_math.py b/test/test_math.py
new file mode 100644
index 0000000000000000000000000000000000000000..c7aabf4687d6d3997d5a37a3c5d4a249143fed1d
--- /dev/null
+++ b/test/test_math.py
@@ -0,0 +1,182 @@
+from __future__ import division
+import math
+import numpy
+import pytools.test
+
+def have_cl():
+ try:
+ import pyopencl
+ return True
+ except:
+ return False
+
+if have_cl():
+ import pyopencl.array as cl_array
+ import pyopencl as cl
+ import pyopencl.clmath as clmath
+ from pyopencl.tools import pytest_generate_tests_for_pyopencl \
+ as pytest_generate_tests
+
+
+
+
+
+sizes = [10, 128, 1024, 1<<10, 1<<13]
+
+
+
+
+def has_double_support(dev):
+ for ext in dev.extensions.split(" "):
+ if ext == "cl_khr_fp64":
+ return True
+ return False
+
+
+
+
+numpy_func_names = {
+ "asin": "arcsin",
+ "acos": "arccos",
+ "atan": "arctan",
+ }
+
+
+
+
+def make_unary_function_test(name, (a, b)=(0, 1), threshold=0):
+ def test(ctx_getter):
+ context = ctx_getter()
+ queue = cl.CommandQueue(context)
+
+ gpu_func = getattr(clmath, name)
+ cpu_func = getattr(numpy, numpy_func_names.get(name, name))
+
+ if has_double_support(context.devices[0]):
+ dtypes = [numpy.float32, numpy.float64]
+ else:
+ dtypes = [numpy.float32]
+
+ for s in sizes:
+ for dtype in dtypes:
+ args = cl_array.arange(context, queue, a, b, (b-a)/s,
+ dtype=numpy.float32)
+ gpu_results = gpu_func(args).get()
+ cpu_results = cpu_func(args.get())
+
+ max_err = numpy.max(numpy.abs(cpu_results - gpu_results))
+ assert (max_err <= threshold).all(), \
+ (max_err, name, dtype)
+
+ return pytools.test.mark_test.opencl(test)
+
+
+
+
+if have_cl():
+ test_ceil = make_unary_function_test("ceil", (-10, 10))
+ test_floor = make_unary_function_test("ceil", (-10, 10))
+ test_fabs = make_unary_function_test("fabs", (-10, 10))
+ test_exp = make_unary_function_test("exp", (-3, 3), 1e-5)
+ test_log = make_unary_function_test("log", (1e-5, 1), 5e-7)
+ test_log10 = make_unary_function_test("log10", (1e-5, 1), 3e-7)
+ test_sqrt = make_unary_function_test("sqrt", (1e-5, 1), 2e-7)
+
+ test_sin = make_unary_function_test("sin", (-10, 10), 1e-7)
+ test_cos = make_unary_function_test("cos", (-10, 10), 1e-7)
+ test_asin = make_unary_function_test("asin", (-0.9, 0.9), 5e-7)
+ test_acos = make_unary_function_test("acos", (-0.9, 0.9), 5e-7)
+ test_tan = make_unary_function_test("tan",
+ (-math.pi/2 + 0.1, math.pi/2 - 0.1), 1e-5)
+ test_atan = make_unary_function_test("atan", (-10, 10), 2e-7)
+
+ test_sinh = make_unary_function_test("sinh", (-3, 3), 1e-6)
+ test_cosh = make_unary_function_test("cosh", (-3, 3), 1e-6)
+ test_tanh = make_unary_function_test("tanh", (-3, 3), 2e-6)
+
+
+
+
+@pytools.test.mark_test.opencl
+def test_fmod(ctx_getter):
+ context = ctx_getter()
+ queue = cl.CommandQueue(context)
+
+ for s in sizes:
+ a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+ a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)/45.2 + 0.1
+ b = clmath.fmod(a, a2)
+
+ a = a.get()
+ a2 = a2.get()
+ b = b.get()
+
+ for i in range(s):
+ assert math.fmod(a[i], a2[i]) == b[i]
+
+@pytools.test.mark_test.opencl
+def test_ldexp(ctx_getter):
+ context = ctx_getter()
+ queue = cl.CommandQueue(context)
+
+ for s in sizes:
+ a = cl_array.arange(context, queue, s, dtype=numpy.float32)
+ a2 = cl_array.arange(context, queue, s, dtype=numpy.float32)*1e-3
+ b = clmath.ldexp(a,a2)
+
+ a = a.get()
+ a2 = a2.get()
+ b = b.get()
+
+ for i in range(s):
+ assert math.ldexp(a[i], int(a2[i])) == b[i]
+
+@pytools.test.mark_test.opencl
+def test_modf(ctx_getter):
+ context = ctx_getter()
+ queue = cl.CommandQueue(context)
+
+ for s in sizes:
+ a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+ fracpart, intpart = clmath.modf(a)
+
+ a = a.get()
+ intpart = intpart.get()
+ fracpart = fracpart.get()
+
+ for i in range(s):
+ fracpart_true, intpart_true = math.modf(a[i])
+
+ assert intpart_true == intpart[i]
+ assert abs(fracpart_true - fracpart[i]) < 1e-4
+
+@pytools.test.mark_test.opencl
+def test_frexp(ctx_getter):
+ context = ctx_getter()
+ queue = cl.CommandQueue(context)
+
+ for s in sizes:
+ a = cl_array.arange(context, queue, s, dtype=numpy.float32)/10
+ significands, exponents = clmath.frexp(a)
+
+ a = a.get()
+ significands = significands.get()
+ exponents = exponents.get()
+
+ for i in range(s):
+ sig_true, ex_true = math.frexp(a[i])
+
+ assert sig_true == significands[i]
+ assert ex_true == exponents[i]
+
+
+
+
+if __name__ == "__main__":
+ # make sure that import failures get reported, instead of skipping the tests.
+ import sys
+ if len(sys.argv) > 1:
+ exec sys.argv[1]
+ else:
+ from py.test.cmdline import main
+ main([__file__])
diff --git a/test/test_wrapper.py b/test/test_wrapper.py
index 56b79369fa159d6ecf5c21507966b823a83a22c2..b61eb4b9bfd44d9d3c304d8eebbd80d43f6a4489 100644
--- a/test/test_wrapper.py
+++ b/test/test_wrapper.py
@@ -17,6 +17,8 @@ def have_cl():
if have_cl():
import pyopencl as cl
+ from pyopencl.tools import pytest_generate_tests_for_pyopencl \
+ as pytest_generate_tests
@@ -182,7 +184,9 @@ class TestCL:
assert not iform.__dict__
@pytools.test.mark_test.opencl
- def test_nonempty_supported_image_formats(self, device, context):
+ def test_nonempty_supported_image_formats(self, device, ctx_getter):
+ context = ctx_getter()
+
if device.image_support:
assert len(cl.get_supported_image_formats(
context, cl.mem_flags.READ_ONLY, cl.mem_object_type.IMAGE2D)) > 0
@@ -191,7 +195,9 @@ class TestCL:
skip("images not supported on %s" % device.name)
@pytools.test.mark_test.opencl
- def test_that_python_args_fail(self, context):
+ def test_that_python_args_fail(self, ctx_getter):
+ context = ctx_getter()
+
prg = cl.Program(context, """
__kernel void mult(__global float *a, float b, int c)
{ a[get_global_id(0)] *= (b+c); }
@@ -220,7 +226,9 @@ class TestCL:
cl.enqueue_read_buffer(queue, a_buf, a_result).wait()
@pytools.test.mark_test.opencl
- def test_image_2d(self, device, context):
+ def test_image_2d(self, device, ctx_getter):
+ context = ctx_getter()
+
if not device.image_support:
from py.test import skip
skip("images not supported on %s" % device)
@@ -270,39 +278,6 @@ class TestCL:
-def pytest_generate_tests(metafunc):
- if have_cl():
- import pyopencl as cl
- else:
- # will still show "skipped" messages
- return
-
- if ("device" in metafunc.funcargnames
- or "context" in metafunc.funcargnames):
- arg_dict = {}
-
- for platform in cl.get_platforms():
- if "platform" in metafunc.funcargnames:
- arg_dict["platform"] = platform
-
- for device in platform.get_devices():
- if "device" in metafunc.funcargnames:
- arg_dict["device"] = device
-
- if "context" in metafunc.funcargnames:
- arg_dict["context"] = cl.Context([device])
-
- metafunc.addcall(funcargs=arg_dict.copy(),
- id=", ".join("%s=%s" % (arg, value)
- for arg, value in arg_dict.iteritems()))
-
- elif "platform" in metafunc.funcargnames:
- for platform in cl.get_platforms():
- metafunc.addcall(
- funcargs=dict(platform=platform),
- id=str(platform))
-
-
if __name__ == "__main__":
# make sure that import failures get reported, instead of skipping the tests.
import pyopencl