compyle
compyle copied to clipboard
Bug: minimum and maximum are computed incorrectly for certain arrays with 'cuda' backend.
This is not an issue for most arrays.
Here is an MWE,
from compyle.array import Array, update_minmax_gpu
import pycuda.gpuarray as gpuarray
from pycuda.reduction import ReductionKernel
import numpy as np
backend = 'cuda'
# numpy
a1 = np.asarray([4.2, 2.0, 0.01, 0.08, 4.0, 29.2], dtype=np.float32)
print(f'{a1=}')
np_max_val = np.max(a1)
np_min_val = np.min(a1)
# compyle
ca1 = Array(dtype=a1.dtype, n=len(a1), backend=backend)
ca1.set(a1)
update_minmax_gpu([ca1], backend=backend)
# cuda
ga1 = gpuarray.to_gpu(a1) # or ca1.dev
max_reduce = ReductionKernel(
np.float32,
neutral=str(np.finfo(a1.dtype).min),
reduce_expr="max(a,b)",
map_expr="x[i]",
arguments="float *x"
)
min_reduce = ReductionKernel(
np.float32,
neutral=str(np.finfo(a1.dtype).max),
reduce_expr="min(a,b)",
map_expr="x[i]",
arguments="float *x"
)
cuda_max_val = max_reduce(ga1).get()
cuda_min_val = min_reduce(ga1).get()
# checks
np.testing.assert_allclose(np_min_val, cuda_min_val)
np.testing.assert_allclose(np_max_val, cuda_max_val)
np.testing.assert_allclose(np_min_val, ca1.minimum)
np.testing.assert_allclose(np_max_val, ca1.maximum)
np.testing.assert_allclose(np_max_val, ca1.maximum)
fails. Here, ca1.maximum
should be 29.2. Instead, it is 4.2.
For some arrays, the minimum is incorrectly computed. I haven't noticed any patterns.
This causes further issues in PySPH's GPUNNPS.
For the same array, backend='opencl' does not give this issue. I will try to find out what's wrong.