diff --git a/CHANGELOG b/CHANGELOG
index 4835105aa..4798ab4bd 100644
--- a/CHANGELOG
+++ b/CHANGELOG
@@ -3,6 +3,7 @@ If not stated, FINUFFT is assumed (cuFINUFFT <=1.3 is listed separately).
 
 Master, using release name V 2.4.0 (2/11/25)
 
+* Added GPU spread interp only test. Added CPU spread interp only test to cmake
 * Make attributes private in Python Plan classes and allow read-only access to
   them using properties (Andén #608).
 * Remove possibility to supply real dtypes to Plan interfaces. Now only complex
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 3736c2caa..68979f3e5 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -8,17 +8,19 @@ set(TESTS
     finufft2dmany_test
     finufft3d_test
     finufft3dmany_test
-    finufft3dkernel_test)
+    finufft3dkernel_test
+    spreadinterp1d_test
+)
 
 foreach(TEST ${TESTS})
-  add_executable(${TEST} ${TEST}.cpp)
-  target_compile_features(${TEST} PRIVATE cxx_std_17)
-  finufft_link_test(${TEST})
-
-  add_executable(${TEST}f ${TEST}.cpp)
-  target_compile_definitions(${TEST}f PRIVATE -DSINGLE)
-  target_compile_features(${TEST}f PRIVATE cxx_std_17)
-  finufft_link_test(${TEST}f)
+    add_executable(${TEST} ${TEST}.cpp)
+    target_compile_features(${TEST} PRIVATE cxx_std_17)
+    finufft_link_test(${TEST})
+
+    add_executable(${TEST}f ${TEST}.cpp)
+    target_compile_definitions(${TEST}f PRIVATE -DSINGLE)
+    target_compile_features(${TEST}f PRIVATE cxx_std_17)
+    finufft_link_test(${TEST}f)
 endforeach()
 
 # copy the DLLs to the build directory so that the tests can find them there fix
@@ -28,75 +30,96 @@ copy_dll(finufft basicpassfail)
 # Add ctest definitions not for both precisions...
 add_executable(testutils testutils.cpp)
 if(FINUFFT_USE_DUCC0)
-  target_compile_definitions(testutils PRIVATE -DFINUFFT_USE_DUCC0)
+    target_compile_definitions(testutils PRIVATE -DFINUFFT_USE_DUCC0)
 endif()
 target_compile_features(testutils PRIVATE cxx_std_17)
 finufft_link_test(testutils)
 add_test(
-  NAME run_testutils
-  COMMAND testutils
-  WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
+    NAME run_testutils
+    COMMAND testutils
+    WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+)
 
 if(NOT FINUFFT_USE_DUCC0)
-  add_executable(fftw_lock_test fftw_lock_test.cpp)
-  target_compile_features(fftw_lock_test PRIVATE cxx_std_17)
-  finufft_link_test(fftw_lock_test)
-
-  add_test(
-    NAME run_fftw_lock_test
-    COMMAND fftw_lock_test
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
+    add_executable(fftw_lock_test fftw_lock_test.cpp)
+    target_compile_features(fftw_lock_test PRIVATE cxx_std_17)
+    finufft_link_test(fftw_lock_test)
+
+    add_test(
+        NAME run_fftw_lock_test
+        COMMAND fftw_lock_test
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
 endif()
 
 # Add ctest definitions that run at both precisions...
 function(add_tests_with_prec PREC REQ_TOL CHECK_TOL SUFFIX)
-  # All of the following should be run at OMP_NUM_THREADS=4 or something small,
-  # as in makefile. This prevents them taking a huge time on a, say, 128-core
-  # Rome node. ... but I don't know platform-indep way to do that! Does anyone?
-
-  add_test(
-    NAME run_basic_pass_fail_${PREC}
-    COMMAND basicpassfail${SUFFIX}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_finufft1d_test_${PREC}
-    COMMAND finufft1d_test${SUFFIX} 1e2 2e2 ${REQ_TOL} 0 2 0.0 ${CHECK_TOL}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_finufft1dmany_test_${PREC}
-    COMMAND finufft1dmany_test${SUFFIX} 3 1e2 2e2 ${REQ_TOL} 0 0 0 2 0.0
+    # All of the following should be run at OMP_NUM_THREADS=4 or something small,
+    # as in makefile. This prevents them taking a huge time on a, say, 128-core
+    # Rome node. ... but I don't know platform-indep way to do that! Does anyone?
+
+    add_test(
+        NAME run_basic_pass_fail_${PREC}
+        COMMAND basicpassfail${SUFFIX}
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_finufft1d_test_${PREC}
+        COMMAND finufft1d_test${SUFFIX} 1e2 2e2 ${REQ_TOL} 0 2 0.0 ${CHECK_TOL}
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_finufft1dmany_test_${PREC}
+        COMMAND
+            finufft1dmany_test${SUFFIX} 3 1e2 2e2 ${REQ_TOL} 0 0 0 2 0.0
             ${CHECK_TOL}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_finufft2d_test_${PREC}
-    COMMAND finufft2d_test${SUFFIX} 1e2 1e1 1e3 ${REQ_TOL} 0 2 0.0 ${CHECK_TOL}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_finufft2dmany_test_${PREC}
-    COMMAND finufft2dmany_test${SUFFIX} 3 1e2 1e1 1e3 ${REQ_TOL} 0 0 0 2 0.0
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_finufft2d_test_${PREC}
+        COMMAND
+            finufft2d_test${SUFFIX} 1e2 1e1 1e3 ${REQ_TOL} 0 2 0.0 ${CHECK_TOL}
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_finufft2dmany_test_${PREC}
+        COMMAND
+            finufft2dmany_test${SUFFIX} 3 1e2 1e1 1e3 ${REQ_TOL} 0 0 0 2 0.0
             ${CHECK_TOL}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_finufft3d_test_${PREC}
-    COMMAND finufft3d_test${SUFFIX} 5 10 20 1e2 ${REQ_TOL} 0 2 0.0 ${CHECK_TOL}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_finufft3dmany_test_${PREC}
-    COMMAND finufft3dmany_test${SUFFIX} 2 5 10 20 1e2 ${REQ_TOL} 0 0 0 2 0.0
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_finufft3d_test_${PREC}
+        COMMAND
+            finufft3d_test${SUFFIX} 5 10 20 1e2 ${REQ_TOL} 0 2 0.0 ${CHECK_TOL}
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_finufft3dmany_test_${PREC}
+        COMMAND
+            finufft3dmany_test${SUFFIX} 2 5 10 20 1e2 ${REQ_TOL} 0 0 0 2 0.0
             ${CHECK_TOL}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
-  add_test(
-    NAME run_dumbinputs_${PREC}
-    COMMAND dumbinputs${SUFFIX}
-    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
-
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME run_dumbinputs_${PREC}
+        COMMAND dumbinputs${SUFFIX}
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
+
+    add_test(
+        NAME spreadinterp1d_${PREC}
+        COMMAND
+            spreadinterp1d_test${SUFFIX} 1e3 1e3 ${REQ_TOL} 0 2 2.0 ${CHECK_TOL}
+        WORKING_DIRECTORY ${CMAKE_BINARY_DIR}
+    )
 endfunction()
 
 # use above function to actually add the tests, with certain requested and check
diff --git a/test/cuda/CMakeLists.txt b/test/cuda/CMakeLists.txt
index cb429c4f4..2aac08711 100644
--- a/test/cuda/CMakeLists.txt
+++ b/test/cuda/CMakeLists.txt
@@ -1,152 +1,258 @@
 file(GLOB test_src "*.c*")
 
 foreach(srcfile ${test_src})
-  string(REGEX REPLACE "\\.c.?" "" executable ${srcfile})
-  get_filename_component(executable ${executable} NAME)
-  add_executable(${executable} ${srcfile})
-  target_include_directories(${executable} PUBLIC ${CUFINUFFT_INCLUDE_DIRS})
-  target_compile_options(${executable}
-                         PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>)
-  find_library(MathLib m)
-  if(MathLib)
-    target_link_libraries(${executable} PUBLIC cufinufft ${MathLib})
-  endif()
-  target_compile_features(${executable} PUBLIC cxx_std_17)
-  set_target_properties(
-    ${executable} PROPERTIES LINKER_LANGUAGE CUDA
-                             CUDA_ARCHITECTURES "${FINUFFT_CUDA_ARCHITECTURES}")
-  message(STATUS "Adding test ${executable}"
-                 " with CUDA_ARCHITECTURES=${FINUFFT_CUDA_ARCHITECTURES}"
-                 " and INCLUDE=${CUFINUFFT_INCLUDE_DIRS}")
+    string(REGEX REPLACE "\\.c.?" "" executable ${srcfile})
+    get_filename_component(executable ${executable} NAME)
+    add_executable(${executable} ${srcfile})
+    target_include_directories(${executable} PUBLIC ${CUFINUFFT_INCLUDE_DIRS})
+    target_compile_options(
+        ${executable}
+        PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>
+    )
+    find_library(MathLib m)
+    if(MathLib)
+        target_link_libraries(${executable} PUBLIC cufinufft ${MathLib})
+    endif()
+    target_compile_features(${executable} PUBLIC cxx_std_17)
+    set_target_properties(
+        ${executable}
+        PROPERTIES
+            LINKER_LANGUAGE CUDA
+            CUDA_ARCHITECTURES "${FINUFFT_CUDA_ARCHITECTURES}"
+    )
+    message(
+        STATUS
+        "Adding test ${executable}"
+        " with CUDA_ARCHITECTURES=${FINUFFT_CUDA_ARCHITECTURES}"
+        " and INCLUDE=${CUFINUFFT_INCLUDE_DIRS}"
+    )
 endforeach()
 
 function(add_tests PREC REQ_TOL CHECK_TOL UPSAMP)
-  add_test(NAME cufinufft1d1_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 0 1 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft1d1_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 1 1 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft1d1_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 2 1 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft1d2_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 0 2 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft1d2_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 1 2 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft1d3_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 0 3 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft1d3_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft1d_test 1 3 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
-                   ${UPSAMP})
-
-  add_test(NAME cufinufft2d1_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 0 1 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d1_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 1 1 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d1_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 2 1 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d2_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 0 2 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d2_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 2 2 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d3_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 0 3 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d3_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2d_test 2 3 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d1many_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2dmany_test 1 1 1e2 2e2 5 0 2e4 ${REQ_TOL}
-                   ${CHECK_TOL} ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d1many_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2dmany_test 2 1 1e2 2e2 5 0 2e4 ${REQ_TOL}
-                   ${CHECK_TOL} ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d2many_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2dmany_test 1 2 1e2 2e2 5 0 2e4 ${REQ_TOL}
-                   ${CHECK_TOL} ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d2many_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2dmany_test 2 2 1e2 2e2 5 0 2e4 ${REQ_TOL}
-                   ${CHECK_TOL} ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d3many_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2dmany_test 1 3 1e2 2e2 5 0 2e4 ${REQ_TOL}
-                   ${CHECK_TOL} ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft2d3many_test_SM_${PREC}_${UPSAMP}
-           COMMAND cufinufft2dmany_test 2 3 1e2 2e2 5 0 2e4 ${REQ_TOL}
-                   ${CHECK_TOL} ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft3d1_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft3d_test 0 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft3d1_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft3d_test 1 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  if(${PREC} STREQUAL "float")
-    add_test(NAME cufinufft3d1_test_SM_${PREC}_${UPSAMP}
-             COMMAND cufinufft3d_test 2 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                     ${PREC} ${UPSAMP})
-
-    add_test(NAME cufinufft3d1_test_block_${PREC}_${UPSAMP}
-             COMMAND cufinufft3d_test 4 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                     ${PREC} ${UPSAMP})
-
-    add_test(NAME cufinufft3d2_test_SM_${PREC}_${UPSAMP}
-             COMMAND cufinufft3d_test 2 2 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                     ${PREC} ${UPSAMP})
-
-    add_test(NAME cufinufft3d3_test_SM_${PREC}_${UPSAMP}
-             COMMAND cufinufft3d_test 2 3 2 5 10 30 ${REQ_TOL} ${CHECK_TOL}
-                     ${PREC} ${UPSAMP})
-  endif()
-
-  add_test(NAME cufinufft3d2_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft3d_test 0 2 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft3d2_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft3d_test 1 2 2 5 10 20 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft3d3_test_auto_${PREC}_${UPSAMP}
-           COMMAND cufinufft3d_test 0 3 2 5 10 30 ${REQ_TOL} ${CHECK_TOL}
-                   ${PREC} ${UPSAMP})
-
-  add_test(NAME cufinufft3d3_test_GM_${PREC}_${UPSAMP}
-           COMMAND cufinufft3d_test 1 3 2 3 7 20 ${REQ_TOL} ${CHECK_TOL}*100
-                   ${PREC} ${UPSAMP})
+    add_test(
+        NAME cufinufft1d1_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 0 1 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft1d1_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 1 1 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft1d1_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 2 1 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft1d2_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 0 2 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft1d2_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 1 2 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft1d3_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 0 3 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft1d3_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft1d_test 1 3 2e3 4e3 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d1_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 0 1 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d1_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 1 1 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d1_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 2 1 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d2_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 0 2 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d2_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 2 2 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d3_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 0 3 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d3_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2d_test 2 3 1e2 2e2 2e4 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d1many_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2dmany_test 1 1 1e2 2e2 5 0 2e4 ${REQ_TOL} ${CHECK_TOL}
+            ${PREC} ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d1many_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2dmany_test 2 1 1e2 2e2 5 0 2e4 ${REQ_TOL} ${CHECK_TOL}
+            ${PREC} ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d2many_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2dmany_test 1 2 1e2 2e2 5 0 2e4 ${REQ_TOL} ${CHECK_TOL}
+            ${PREC} ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d2many_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2dmany_test 2 2 1e2 2e2 5 0 2e4 ${REQ_TOL} ${CHECK_TOL}
+            ${PREC} ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d3many_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2dmany_test 1 3 1e2 2e2 5 0 2e4 ${REQ_TOL} ${CHECK_TOL}
+            ${PREC} ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft2d3many_test_SM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft2dmany_test 2 3 1e2 2e2 5 0 2e4 ${REQ_TOL} ${CHECK_TOL}
+            ${PREC} ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft3d1_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft3d_test 0 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft3d1_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft3d_test 1 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    if(${PREC} STREQUAL "float")
+        add_test(
+            NAME cufinufft3d1_test_SM_${PREC}_${UPSAMP}
+            COMMAND
+                cufinufft3d_test 2 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+                ${UPSAMP}
+        )
+
+        add_test(
+            NAME cufinufft3d1_test_block_${PREC}_${UPSAMP}
+            COMMAND
+                cufinufft3d_test 4 1 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+                ${UPSAMP}
+        )
+
+        add_test(
+            NAME cufinufft3d2_test_SM_${PREC}_${UPSAMP}
+            COMMAND
+                cufinufft3d_test 2 2 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+                ${UPSAMP}
+        )
+
+        add_test(
+            NAME cufinufft3d3_test_SM_${PREC}_${UPSAMP}
+            COMMAND
+                cufinufft3d_test 2 3 2 5 10 30 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+                ${UPSAMP}
+        )
+    endif()
+
+    add_test(
+        NAME cufinufft3d2_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft3d_test 0 2 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft3d2_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft3d_test 1 2 2 5 10 20 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft3d3_test_auto_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft3d_test 0 3 2 5 10 30 ${REQ_TOL} ${CHECK_TOL} ${PREC}
+            ${UPSAMP}
+    )
+
+    add_test(
+        NAME cufinufft3d3_test_GM_${PREC}_${UPSAMP}
+        COMMAND
+            cufinufft3d_test 1 3 2 3 7 20 ${REQ_TOL} ${CHECK_TOL}*100 ${PREC}
+            ${UPSAMP}
+    )
 endfunction()
 
 add_test(NAME cufinufft_public_api COMMAND public_api_test)
 add_test(NAME cufinufft_makeplan COMMAND test_makeplan)
 add_test(NAME cufinufft_math_test COMMAND cufinufft_math_test)
+add_test(
+    NAME cufinufft1dspreadinterponly_double_test
+    COMMAND cufinufft1dspreadinterponly_test 1e3 1e6 1e-6 1e-5 d 2.00
+)
+add_test(
+    NAME cufinufft1dspreadinterponly_float_test
+    COMMAND cufinufft1dspreadinterponly_test 1e3 1e6 1e-5 1e-4 f 1.25
+)
 
 add_tests(float 1e-5 2e-4 2.0)
 add_tests(double 1e-12 1e-11 2.0)
@@ -155,7 +261,7 @@ add_tests(double 1e-8 1e-7 1.25)
 # the upsamp is appended to the testname, ctest does not allows multiple tests
 # to share the same testname hence we use the trick 0. and 0.f to differentiate
 # the tests and allow them to run in the future we should add the precision to
-# the test (f
+# the test
 add_tests(float 1e-5 2e-4 0.f)
 add_tests(double 1e-12 1e-11 0.f)
 add_tests(float 1e-5 2e-4 0.)
diff --git a/test/cuda/cufinufft1dspreadinterponly_test.cu b/test/cuda/cufinufft1dspreadinterponly_test.cu
new file mode 100644
index 000000000..9b7f80b6f
--- /dev/null
+++ b/test/cuda/cufinufft1dspreadinterponly_test.cu
@@ -0,0 +1,222 @@
+#include <cmath>
+#include <complex>
+#include <iomanip>
+#include <iostream>
+#include <random>
+
+#include <cufinufft.h>
+
+#include <cufinufft/contrib/helper_cuda.h>
+#include <cufinufft/impl.h>
+#include <cufinufft/utils.h>
+
+#include <thrust/complex.h>
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+
+using cufinufft::utils::infnorm;
+
+template<typename T>
+int run_test(int N1, int M, T tol, T checktol, int iflag, double upsampfac) {
+  // tol and upsamplefac are used to determine the kernel
+
+  std::cout << std::scientific << std::setprecision(3);
+  int ier{};
+
+  const int dim = 1;
+
+  // Here we setup our own opts, for gpu_method.
+  cufinufft_opts opts;
+  cufinufft_default_opts(&opts);
+
+  // opts.gpu_method       = method;
+  opts.gpu_maxbatchsize     = 1;
+  opts.gpu_spreadinterponly = 1;
+  opts.upsampfac            = upsampfac;
+
+  int ntransf   = 1;
+  int nmodes[3] = {N1, 1, 1};
+
+  cufinufft_plan_t<T> *dplan;
+
+  thrust::host_vector<T> x(M);
+  thrust::host_vector<thrust::complex<T>> c(M);
+  thrust::host_vector<thrust::complex<T>> fk(M);
+
+  thrust::device_vector<T> d_x(M);
+  thrust::device_vector<thrust::complex<T>> d_c(M);
+  thrust::device_vector<thrust::complex<T>> d_fk(M);
+
+  x[0] = 0.0;
+  c[0] = 1.0;
+  d_x  = x;
+  d_c  = c;
+
+  cudaEvent_t start, stop;
+  float milliseconds = 0;
+
+  cudaEventCreate(&start);
+  cudaEventCreate(&stop);
+
+  ier = cufinufft_makeplan_impl<T>(1, dim, nmodes, iflag, ntransf, tol, &dplan, &opts);
+  if (ier != 0) {
+    printf("err: cufinufft1d_plan (ier=%d)\n", ier);
+    return ier;
+  }
+  ier = cufinufft_setpts_impl<T>(M, d_x.data().get(), nullptr, nullptr, 0, nullptr,
+                                 nullptr, nullptr, dplan);
+  if (ier != 0) {
+    printf("err: cufinufft_setpts (ier=%d)\n", ier);
+    return ier;
+  }
+  ier = cufinufft_execute_impl<T>((cuda_complex<T> *)d_c.data().get(),
+                                  (cuda_complex<T> *)d_fk.data().get(), dplan);
+  if (ier != 0) {
+    printf("err: cufinufft1d_exec (ier=%d)\n", ier);
+    return ier;
+  }
+  cufinufft_destroy_impl(dplan);
+
+  fk = d_fk;
+  const auto kersum =
+      std::accumulate(fk.begin(), fk.end(), thrust::complex<T>(T(0), T(0)));
+
+  // making data
+  std::default_random_engine eng(1);
+  std::uniform_real_distribution<T> dist11(-1, 1);
+  auto randm11 = [&eng, &dist11]() {
+    return dist11(eng);
+  };
+
+  // Making data
+  for (int i = 0; i < M; i++) {
+    x[i] = M_PI * randm11(); // x in [-pi,pi)
+  }
+
+  for (int i = 0; i < M; i++) {
+    c[i].real(randm11());
+    c[i].imag(randm11());
+  }
+
+  d_x = x;
+  d_c = c;
+
+  printf("spread-only test 1d:\n"); // ............................................
+
+  cudaDeviceSynchronize();
+  cudaEventRecord(start);
+
+  ier = cufinufft_makeplan_impl<T>(1, dim, nmodes, iflag, ntransf, tol, &dplan, &opts);
+  if (ier != 0) {
+    printf("err: cufinufft1d_plan (ier=%d)\n", ier);
+    return ier;
+  }
+  ier = cufinufft_setpts_impl<T>(M, d_x.data().get(), nullptr, nullptr, 0, nullptr,
+                                 nullptr, nullptr, dplan);
+  if (ier != 0) {
+    printf("err: cufinufft_setpts (ier=%d)\n", ier);
+    return ier;
+  }
+  ier = cufinufft_execute_impl<T>((cuda_complex<T> *)d_c.data().get(),
+                                  (cuda_complex<T> *)d_fk.data().get(), dplan);
+
+  if (ier != 0) {
+    printf("err: cufinufft1d_exec (ier=%d)\n", ier);
+    return ier;
+  }
+  cufinufft_destroy_impl(dplan);
+
+  cudaEventRecord(stop);
+  cudaEventSynchronize(stop);
+  cudaEventElapsedTime(&milliseconds, start, stop);
+
+  printf("\t%lld NU pts spread to %lld grid in %.3g s \t%.3g NU pts/s\n",
+         static_cast<long long>(M), static_cast<long long>(N1), milliseconds / 1000,
+         T(M) / (milliseconds / 1000));
+
+  fk = d_fk;
+
+  auto csum       = std::accumulate(c.begin(), c.end(), thrust::complex<T>(T(0), T(0)));
+  const auto mass = std::accumulate(fk.begin(), fk.end(), thrust::complex<T>(T(0), T(0)));
+  const auto rel_mass_err = thrust::abs(mass - kersum * csum) / thrust::abs(mass);
+  printf("\trel mass err %.3g\n", rel_mass_err);
+
+  printf("interp-only test 1d:\n"); // ............................................
+
+  std::fill(fk.begin(), fk.end(), thrust::complex<T>(T(1), T(0)));
+  d_fk = fk;
+
+  cudaDeviceSynchronize();
+  cudaEventRecord(start);
+
+  ier = cufinufft_makeplan_impl<T>(2, dim, nmodes, iflag, ntransf, tol, &dplan, &opts);
+  if (ier != 0) {
+    printf("err: cufinufft1d_plan (ier=%d)\n", ier);
+    return ier;
+  }
+  ier = cufinufft_setpts_impl<T>(M, d_x.data().get(), nullptr, nullptr, 0, nullptr,
+                                 nullptr, nullptr, dplan);
+  if (ier != 0) {
+    printf("err: cufinufft_setpts (ier=%d)\n", ier);
+    return ier;
+  }
+
+  ier = cufinufft_execute_impl<T>((cuda_complex<T> *)d_c.data().get(),
+                                  (cuda_complex<T> *)d_fk.data().get(), dplan);
+
+  if (ier != 0) {
+    printf("err: cufinufft1d_exec (ier=%d)\n", ier);
+    return ier;
+  }
+  cufinufft_destroy_impl(dplan);
+
+  cudaEventSynchronize(stop);
+  cudaEventRecord(stop);
+  cudaEventElapsedTime(&milliseconds, start, stop);
+
+  printf("\t%lld NU pts interp from %lld grid in %.3g s \t%.3g NU pts/s\n",
+         static_cast<long long>(M), static_cast<long long>(N1), milliseconds / 1000,
+         T(M) / (milliseconds / 1000));
+
+  c = d_c;
+
+  csum         = std::accumulate(c.begin(), c.end(), thrust::complex<T>(T(0), T(0)));
+  auto sup_err = T(0.0);
+  for (auto cj : c) sup_err = std::max(sup_err, abs(cj - kersum));
+  const auto rel_sup_err = sup_err / thrust::abs(kersum);
+  printf("\trel sup err %.3g\n", rel_sup_err);
+
+  cudaEventDestroy(start);
+  cudaEventDestroy(stop);
+
+  const auto rel_error = std::max(rel_mass_err, rel_sup_err);
+  return std::isnan(rel_error) || rel_error > checktol;
+}
+
+int main(int argc, char *argv[]) {
+  if (argc != 7) {
+    fprintf(stderr,
+            "Usage: cufinufft1dspreadinterponly_test N1 M tol checktol prec upsampfac\n"
+            "Arguments:\n"
+            "  N1: Number of fourier modes\n"
+            "  M: The number of non-uniform points\n"
+            "  tol: NUFFT tolerance\n"
+            "  checktol:  relative error to pass test\n"
+            "  precision: f or d\n"
+            "  upsampfac: upsampling factor\n");
+    return 1;
+  }
+  const int N1           = atof(argv[1]);
+  const int M            = atof(argv[2]);
+  const double tol       = atof(argv[3]);
+  const double checktol  = atof(argv[4]);
+  const int iflag        = 1;
+  const char prec        = argv[5][0];
+  const double upsampfac = atof(argv[6]);
+  if (prec == 'f')
+    return run_test<float>(N1, M, tol, checktol, iflag, upsampfac);
+  else if (prec == 'd')
+    return run_test<double>(N1, M, tol, checktol, iflag, upsampfac);
+  else
+    return -1;
+}