From 9496eeaac33e86a7a11d9f33b5a6a2d0e365818e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Joakim=20And=C3=A9n?= Date: Fri, 31 May 2024 04:45:26 -0400 Subject: [PATCH 1/6] 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). --- include/cufinufft.h | 8 ++++---- src/cuda/cufinufft.cu | 16 ++++++++++------ 2 files changed, 14 insertions(+), 10 deletions(-) diff --git a/include/cufinufft.h b/include/cufinufft.h index b323d94c0..ebd2798fc 100644 --- a/include/cufinufft.h +++ b/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, const 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, const 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); diff --git a/src/cuda/cufinufft.cu b/src/cuda/cufinufft.cu index 40510e95b..26b64c47e 100644 --- a/src/cuda/cufinufft.cu +++ b/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 **)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::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 *)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::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 *)d_plan); } From 96f450f38e6957e3b90c15959afa903258a3cd44 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Joakim=20And=C3=A9n?= Date: Thu, 28 Dec 2023 00:43:54 +0100 Subject: [PATCH 2/6] cuda+py: update py interface to deal with 64-bit M --- python/cufinufft/cufinufft/_cufinufft.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/cufinufft/cufinufft/_cufinufft.py b/python/cufinufft/cufinufft/_cufinufft.py index 2fccba192..bbd446a60 100644 --- a/python/cufinufft/cufinufft/_cufinufft.py +++ b/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 From 6f30fbda4fbff21c270575abc27513e894356b8a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Joakim=20And=C3=A9n?= Date: Fri, 31 May 2024 04:47:16 -0400 Subject: [PATCH 3/6] cuda: update CHANGELOG wrt new setpts API --- CHANGELOG | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/CHANGELOG b/CHANGELOG index d89a2f1ac..dabf49c8b 100644 --- a/CHANGELOG +++ b/CHANGELOG @@ -22,6 +22,14 @@ If not stated, FINUFFT is assumed (cuFINUFFT <=1.3 is listed separately). 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 listed in the future. + +>>>>>>> 3d9f5a2e (cuda: update CHANGELOG wrt new setpts API) V 2.2.0 (12/12/23) From 7d740ef4a9d062072f2b9a4b67d728e89bb46cf4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Joakim=20And=C3=A9n?= Date: Tue, 11 Jun 2024 09:02:40 +0200 Subject: [PATCH 4/6] ci: no cache for Jenkins --- Jenkinsfile | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index bcb60c9c8..6600c1cc3 100644 --- a/Jenkinsfile +++ b/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 From 3c0ab9aee07c6b57e9d7d2a2fd788d92ed22352e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Joakim=20And=C3=A9n?= Date: Tue, 25 Jun 2024 19:47:36 +0200 Subject: [PATCH 5/6] docs: fix typos in CHANGELOG --- CHANGELOG | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/CHANGELOG b/CHANGELOG index dabf49c8b..d35999aa2 100644 --- a/CHANGELOG +++ b/CHANGELOG @@ -27,9 +27,7 @@ as opposed to 32-bit. While this does modify the ABI, most code will just need t 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 listed in the future. - ->>>>>>> 3d9f5a2e (cuda: update CHANGELOG wrt new setpts API) +This restriction may be lifted in the future. V 2.2.0 (12/12/23) From 0ce56421dbe2258675b8612854861d389096d8f1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Joakim=20And=C3=A9n?= Date: Tue, 25 Jun 2024 19:49:04 +0200 Subject: [PATCH 6/6] cuda: remove const qualifiers for setpts --- include/cufinufft.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/cufinufft.h b/include/cufinufft.h index ebd2798fc..52dd781a0 100644 --- a/include/cufinufft.h +++ b/include/cufinufft.h @@ -19,9 +19,9 @@ 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, const int64_t M, double *d_x, double *d_y, +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, const int64_t M, float *d_x, float *d_y, +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);