Bug: minimum and maximum are computed incorrectly for certain arrays with 'cuda' backend.

[nauaneed]

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.