ENH: cupyx/scipy/interpolate: add griddata

[ev-br]

Port griddata function from scipy.interpolate. This is a pure python convenience wrapper ofLinearNDInterpolator, NearestNDInterpolator and CloughTocher2DInterpolator, so no need to benchmark performance.

This PR is on top of #8220 for the NearestNDInterpolator and #8289 for interp1d. Will need a rebase or two when these two PRs land.

[mergify]

This pull request is now in conflicts. Could you fix it @ev-br? 🙏

[mergify]

This pull request is now in conflicts. Could you fix it @ev-br? 🙏

[ev-br]

Both dependencies are merged, the PR is rebased on main and is ready for review.

[takagi]

/test mini

[takagi]

LGTM! Waiting for CI pass.

[takagi]

Seems that CI fails with illegal memory access. Does this occurs on your local? Disabling CuPy's memory pool cupy.cuda.set_allocator(None) may help.

E   cupy_backends.cuda.api.driver.CUDADriverError: CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered

[ev-br]

Sadly, no, I don't see this locally. $ pytest tests/cupyx_tests/scipy_tests/interpolate_tests/ -v passes for me (on CUDA 12 and scipy 1.13) modulo a test failure due to a unrelated scipy 1.13 change, scipy/scipy#20364.

Also given that this PR only adds a thin pure python wrapper, the the root problem is likely somewhere else. Where does the error originate?

[takagi]

Seems that its first occurrence in CI log is TestGriddata.test_fill_value , but as you say it may not the root problem. I will also check this issue.

https://ci.preferred.jp/cupy.linux.cuda112/161482/
https://storage.googleapis.com/chainer-artifacts-pfn-public-ci/cupy-ci/161482/log.txt

...
cupyx_tests/scipy_tests/interpolate_tests/test_ndgriddata.py ..sssFFFFFF [ 36%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF                                      [ 36%]
cupyx_tests/scipy_tests/interpolate_tests/test_polyint.py FFFFFFFFFFFFFF [ 36%]
sFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFssssssssssssssssssssssssssssssssssssssss [ 37%]
ssssssssssssssssssssFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFsFFFFFFFFFFFFFsFFF [ 37%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF [ 37%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF [ 37%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFssssssssFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF [ 37%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFsssss                                        [ 37%]
cupyx_tests/scipy_tests/interpolate_tests/test_ppoly.py FFFFFFFFFFFFFFFF [ 37%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFssssssssFFFFFFFFFFFFsssFFFFFFFFFFFFFFF [ 37%]
FFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFF                                 [ 37%]
cupyx_tests/scipy_tests/interpolate_tests/test_rbfinterp.py ssssssssFFFF [ 37%]
FFFFFFF

=================================== FAILURES ===================================
_________________________ TestGriddata.test_fill_value _________________________

self = <cupyx_tests.scipy_tests.interpolate_tests.test_ndgriddata.TestGriddata object at 0x7fb8d45a5df0>
xp = <module 'cupy' from '/root/.local/lib/python3.9/site-packages/cupy/__init__.py'>
scp = <module 'cupyx.scipy' from '/root/.local/lib/python3.9/site-packages/cupyx/scipy/__init__.py'>

    @testing.numpy_cupy_allclose(scipy_name='scp')
    def test_fill_value(self, xp, scp):
        x = [(0, 0), (0, 1), (1, 0)]
        y = [1, 2, 3]
...

[ev-br]

TBH, I don't see how a pure python addition of this PR can trigger this. The only possiblility is that a nan assignment is broken somehow, in this line: https://github.com/cupy/cupy/pull/8290/files#diff-e6ca1b203fa4c1b7e7f6170aa16844edb19aa250d65f43d99d4bd3ee26187b36R256 with the default of fill_value=cupy.nan. But this is definitely not the only place nans appear, so it's unlikely too.

[kmaehashi]

I've tried running compute-sanitizer --tool memcheck python -m pytest -v -s -k test_fill_value tests/cupyx_tests/scipy_tests/interpolate_tests/test_ndgriddata.py on my local with memory pool disabled by commenting out these lines. It's not always reproducing (it even sometimes hangs) but I got this error reported:

========= Invalid __global__ read of size 4 bytes
=========     at kerFindClosestTri(Point2 *, int, Tri *, int, TriOpp *, long long *, unsigned int *, Point2 *, Point2 *, const double *, const double *, Point2 *, Point2 *, double, bool, int *, double *)+0x85d0
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7f2ec2a007f4 is out of bounds
=========     and is 12 bytes before the nearest allocation at 0x7f2ec2a00800 of size 48 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x32e94f]

[takagi]

@andfoy Can I ask you to check if kerFindClosestTri could make any out-of-bounds access again?

[andfoy]

Thanks for the report @takagi, the reported segfault should be fixed by #8363