cuda: update interface to take 64-bit M (#411)

* cuda: change to int64_t for number of pts

While the interface changes, we still won't allow more than 2e9
points (for now) since this will complicate the code quite a bit with
little tangible benefits (transforms these size are currently out of
range for most GPUs in terms of memory consumption).

* cuda+py: update py interface to deal with 64-bit M

* cuda: update CHANGELOG wrt new setpts API

* ci: no cache for Jenkins

* docs: fix typos in CHANGELOG

* cuda: remove const qualifiers for setpts

Author: janden

CHANGELOG

@@ -29,6 +29,12 @@ V 2.3.0beta (6/21/24)
   Created a .clang-format file to define the style similar to the existing style.
   Applied clang-format to all cmake, C, C++, and CUDA code. Ignored the blame
   using .git-blame-ignore-revs. Added a contributing.md for developers.
+* cuFINUFFT interface update: number of nonuniform points M is now a 64-bit integer
+as opposed to 32-bit. While this does modify the ABI, most code will just need to
+recompile against the new library as compilers will silently upcast any 32-bit
+integers to 64-bit when calling cufinufft(f)_setpts. Note that internally, 32-bit
+integers are still used, so calling cufinufft with more than 2e9 points will fail.
+This restriction may be lifted in the future.
 
 V 2.2.0 (12/12/23)
 

Jenkinsfile

@@ -46,11 +46,11 @@ pipeline {
     sh '${PYBIN}/python3 -m venv $HOME'
     sh '''#!/bin/bash -ex
       source $HOME/bin/activate
-      python3 -m pip install --upgrade pip
-      python3 -m pip install --upgrade pycuda cupy-cuda112 numba
-      python3 -m pip install torch==1.10.2+cu111 -f https://download.pytorch.org/whl/torch_stable.html
-      python3 -m pip install python/cufinufft
-      python3 -m pip install pytest
+      python3 -m pip install --no-cache-dir --upgrade pip
+      python3 -m pip install --no-cache-dir --upgrade pycuda cupy-cuda112 numba
+      python3 -m pip install --no-cache-dir torch==1.10.2+cu111 -f https://download.pytorch.org/whl/torch_stable.html
+      python3 -m pip install --no-cache-dir python/cufinufft
+      python3 -m pip install --no-cache-dir pytest
       python -c "from numba import cuda; cuda.cudadrv.libs.test()"
       python3 -m pytest --framework=pycuda python/cufinufft
       python3 -m pytest --framework=numba python/cufinufft

include/cufinufft.h

@@ -19,10 +19,10 @@ int cufinufft_makeplan(int type, int dim, const int64_t *n_modes, int iflag, int
 int cufinufftf_makeplan(int type, int dim, const int64_t *n_modes, int iflag, int ntr,
                         float eps, cufinufftf_plan *d_plan_ptr, cufinufft_opts *opts);
 
-int cufinufft_setpts(cufinufft_plan d_plan, int M, double *d_x, double *d_y, double *d_z,
-                     int N, double *d_s, double *d_t, double *d_u);
-int cufinufftf_setpts(cufinufftf_plan d_plan, int M, float *d_x, float *d_y, float *d_z,
-                      int N, float *d_s, float *d_t, float *d_u);
+int cufinufft_setpts(cufinufft_plan d_plan, int64_t M, double *d_x, double *d_y,
+                     double *d_z, int N, double *d_s, double *d_t, double *d_u);
+int cufinufftf_setpts(cufinufftf_plan d_plan, int64_t M, float *d_x, float *d_y,
+                      float *d_z, int N, float *d_s, float *d_t, float *d_u);
 
 int cufinufft_execute(cufinufft_plan d_plan, cuDoubleComplex *d_c, cuDoubleComplex *d_fk);
 int cufinufftf_execute(cufinufftf_plan d_plan, cuFloatComplex *d_c, cuFloatComplex *d_fk);

python/cufinufft/cufinufft/_cufinufft.py

@@ -102,13 +102,13 @@ class NufftOpts(ctypes.Structure):
 
 _set_pts = lib.cufinufft_setpts
 _set_pts.argtypes = [
-    c_void_p, c_int, c_void_p, c_void_p, c_void_p, ctypes.c_int, c_double_p,
+    c_void_p, c_int64, c_void_p, c_void_p, c_void_p, ctypes.c_int, c_double_p,
     c_double_p, c_double_p]
 _set_pts.restype = c_int
 
 _set_ptsf = lib.cufinufftf_setpts
 _set_ptsf.argtypes = [
-    c_void_p, c_int, c_void_p, c_void_p, c_void_p, ctypes.c_int, c_float_p,
+    c_void_p, c_int64, c_void_p, c_void_p, c_void_p, ctypes.c_int, c_float_p,
     c_float_p, c_float_p]
 _set_ptsf.restype = c_int
 

src/cuda/cufinufft.cu

@@ -48,15 +48,19 @@ int cufinufft_makeplan(int type, int dim, const int64_t *nmodes, int iflag, int
                                  (cufinufft_plan_t<double> **)d_plan_ptr, opts);
 }
 
-int cufinufftf_setpts(cufinufftf_plan d_plan, int M, float *d_x, float *d_y, float *d_z,
-                      int N, float *d_s, float *d_t, float *d_u) {
-  return cufinufft_setpts_impl(M, d_x, d_y, d_z, N, d_s, d_t, d_u,
+int cufinufftf_setpts(cufinufftf_plan d_plan, const int64_t M, float *d_x, float *d_y,
+                      float *d_z, int N, float *d_s, float *d_t, float *d_u) {
+  if (M > std::numeric_limits<int32_t>::max()) return FINUFFT_ERR_NDATA_NOTVALID;
+
+  return cufinufft_setpts_impl((int)M, d_x, d_y, d_z, N, d_s, d_t, d_u,
                                (cufinufft_plan_t<float> *)d_plan);
 }
 
-int cufinufft_setpts(cufinufft_plan d_plan, int M, double *d_x, double *d_y, double *d_z,
-                     int N, double *d_s, double *d_t, double *d_u) {
-  return cufinufft_setpts_impl(M, d_x, d_y, d_z, N, d_s, d_t, d_u,
+int cufinufft_setpts(cufinufft_plan d_plan, const int64_t M, double *d_x, double *d_y,
+                     double *d_z, int N, double *d_s, double *d_t, double *d_u) {
+  if (M > std::numeric_limits<int32_t>::max()) return FINUFFT_ERR_NDATA_NOTVALID;
+
+  return cufinufft_setpts_impl((int)M, d_x, d_y, d_z, N, d_s, d_t, d_u,
                                (cufinufft_plan_t<double> *)d_plan);
 }