diff --git a/CMakeLists.txt b/CMakeLists.txt
index d1f685a..e6bbf15 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -13,8 +13,13 @@ if (NOT CMAKE_BUILD_TYPE)
 endif()
 
 # User-defined build options
+option(CUDECOMP_BUILD_CPU_ONLY "Build the toolkit-free, MPI-only CPU debug/test variant" OFF)
+set(CUDECOMP_NVTX_DEFAULT ON)
+if (CUDECOMP_BUILD_CPU_ONLY)
+  set(CUDECOMP_NVTX_DEFAULT OFF)
+endif()
 option(CUDECOMP_BUILD_FORTRAN "Build Fortran bindings" ON)
-option(CUDECOMP_ENABLE_NVTX "Enable NVTX ranges" ON)
+option(CUDECOMP_ENABLE_NVTX "Enable NVTX ranges" ${CUDECOMP_NVTX_DEFAULT})
 option(CUDECOMP_ENABLE_NVSHMEM "Enable NVSHMEM" OFF)
 option(CUDECOMP_BUILD_TESTS "Build tests" OFF)
 option(CUDECOMP_TEST_FETCH_GTEST "Fetch GoogleTest if a system installation is unavailable" OFF)
@@ -22,23 +27,31 @@ option(CUDECOMP_BUILD_EXTRAS "Build benchmark, examples, and tests" OFF)
 set(CUDECOMP_NCCL_HOME CACHE STRING "Path to search for NCCL installation. Use to override NVHPC provided NCCL version.")
 set(CUDECOMP_NVSHMEM_HOME CACHE STRING "Path to search for NVSHMEM installation. Use to override NVHPC provided NVSHMEM version.")
 
-# Use NVHPC compilers by default
-find_program(NVHPC_CXX_BIN "nvc++" REQUIRED)
-set(CMAKE_CXX_COMPILER ${NVHPC_CXX_BIN})
-
-if (CUDECOMP_BUILD_FORTRAN)
-  find_program(NVHPC_Fortran_BIN "nvfortran" REQUIRED)
-  set(CMAKE_Fortran_COMPILER ${NVHPC_Fortran_BIN})
+if (CUDECOMP_BUILD_CPU_ONLY AND (CUDECOMP_ENABLE_NVTX OR CUDECOMP_ENABLE_NVSHMEM))
+  message(FATAL_ERROR "CPU-only builds do not support NVTX or NVSHMEM")
 endif()
 
-# Locate and use NVHPC CMake configuration
-string(REPLACE "compilers/bin/nvc++" "cmake" NVHPC_CMAKE_DIR ${NVHPC_CXX_BIN})
-set(CMAKE_PREFIX_PATH ${NVHPC_CMAKE_DIR})
+if (NOT CUDECOMP_BUILD_CPU_ONLY)
+  # Use NVHPC compilers by default
+  find_program(NVHPC_CXX_BIN "nvc++" REQUIRED)
+  set(CMAKE_CXX_COMPILER ${NVHPC_CXX_BIN})
 
+  if (CUDECOMP_BUILD_FORTRAN)
+    find_program(NVHPC_Fortran_BIN "nvfortran" REQUIRED)
+    set(CMAKE_Fortran_COMPILER ${NVHPC_Fortran_BIN})
+  endif()
+
+  # Locate and use NVHPC CMake configuration
+  string(REPLACE "compilers/bin/nvc++" "cmake" NVHPC_CMAKE_DIR ${NVHPC_CXX_BIN})
+  set(CMAKE_PREFIX_PATH ${NVHPC_CMAKE_DIR})
+endif()
+
+set(LANGS CXX)
+if (NOT CUDECOMP_BUILD_CPU_ONLY)
+  list(APPEND LANGS CUDA)
+endif()
 if (CUDECOMP_BUILD_FORTRAN)
-  set(LANGS CXX CUDA Fortran)
-else()
-  set(LANGS CXX CUDA)
+  list(APPEND LANGS Fortran)
 endif()
 
 set(CUDECOMP_VERSION_COMPONENTS)
@@ -58,33 +71,36 @@ set(CUDECOMP_SOVERSION 0)
 
 project(cudecomp VERSION ${CUDECOMP_VERSION} LANGUAGES ${LANGS})
 
-if (CUDECOMP_BUILD_FORTRAN)
+if (CUDECOMP_BUILD_FORTRAN AND NOT CUDECOMP_BUILD_CPU_ONLY)
   # Remove debug information to enable mixed Fortran compilation with cc100 and older compute capabilities.
   string(REGEX REPLACE "(^| )-g(opt)?($| )" "\\3" CMAKE_Fortran_FLAGS_RELWITHDEBINFO
                        "${CMAKE_Fortran_FLAGS_RELWITHDEBINFO}")
 endif()
 
-# Set up CUDA compute capabilities by CUDA version. Users can override defaults with CUDECOMP_CUDA_CC_LIST
-if (NOT CUDECOMP_CUDA_CC_LIST)
-  if (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "75;80;90;100")
-  elseif (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8)
-    set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "70;80;90;100")
-  elseif (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 11.8)
-    set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "70;80;90")
-  else()
-    set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "70;80")
+if (NOT CUDECOMP_BUILD_CPU_ONLY)
+  # Set up CUDA compute capabilities by CUDA version. Users can override defaults with CUDECOMP_CUDA_CC_LIST
+  if (NOT CUDECOMP_CUDA_CC_LIST)
+    if (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+      set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "75;80;90;100")
+    elseif (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8)
+      set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "70;80;90;100")
+    elseif (${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 11.8)
+      set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "70;80;90")
+    else()
+      set(CUDECOMP_CUDA_CC_LIST_DEFAULTS "70;80")
+    endif()
+
+    set(CUDECOMP_CUDA_CC_LIST ${CUDECOMP_CUDA_CC_LIST_DEFAULTS} CACHE STRING "List of CUDA compute capabilities to build cuDecomp for.")
   endif()
 
-  set(CUDECOMP_CUDA_CC_LIST ${CUDECOMP_CUDA_CC_LIST_DEFAULTS} CACHE STRING "List of CUDA compute capabilities to build cuDecomp for.")
-endif()
+  # Detect if Cray compiler wrappers are available to assess if in Cray environment.
+  # We do not use the Cray compiler wrappers directly for greater flexibility.
+  find_program(CRAY_CC_BIN "CC")
 
-# Detect if Cray compiler wrappers are available to assess if in Cray environment.
-# We do not use the Cray compiler wrappers directly for greater flexibility.
-find_program(CRAY_CC_BIN "CC")
+  if (CRAY_CC_BIN)
+    message(STATUS "Found Cray CC wrapper. Compiling for Cray programming environment.")
+  endif()
 
-if (CRAY_CC_BIN)
-  message(STATUS "Found Cray CC wrapper. Compiling for Cray programming environment.")
 endif()
 
 # MPI
@@ -94,74 +110,77 @@ if (CUDECOMP_BUILD_TESTS)
 endif()
 find_package(MPI REQUIRED)
 
-if (CRAY_CC_BIN)
-  # FindMPI does not include Cray GTL (e.g. CUDA-aware) libs
-  # automatically in Cray environment. Locate it to include in linking.
-  string(REPLACE ":" ";" CRAY_LIB_PATHS $ENV{CRAY_LD_LIBRARY_PATH})
-  find_library(CRAY_MPI_GTL_CUDA_LIBRARY REQUIRED
-    NAMES mpi_gtl_cuda
-    HINTS ${CRAY_LIB_PATHS}
-  )
-
-  # Cray GTL libs benefit from linking against gdrcopy, so also
-  # locating that library.
-  find_library(GDRCOPY_LIBRARY REQUIRED
-    NAMES gdrapi
-  )
-
-  message(STATUS "Found Cray GTL library: " ${CRAY_MPI_GTL_CUDA_LIBRARY})
-  message(STATUS "Found GDRCopy library: " ${GDRCOPY_LIBRARY})
-endif()
+if (NOT CUDECOMP_BUILD_CPU_ONLY)
+  if (CRAY_CC_BIN)
+    # FindMPI does not include Cray GTL (e.g. CUDA-aware) libs
+    # automatically in Cray environment. Locate it to include in linking.
+    string(REPLACE ":" ";" CRAY_LIB_PATHS $ENV{CRAY_LD_LIBRARY_PATH})
+    find_library(CRAY_MPI_GTL_CUDA_LIBRARY REQUIRED
+      NAMES mpi_gtl_cuda
+      HINTS ${CRAY_LIB_PATHS}
+    )
 
-# HPC SDK
-find_package(NVHPC REQUIRED COMPONENTS CUDA MATH)
+    # Cray GTL libs benefit from linking against gdrcopy, so also
+    # locating that library.
+    find_library(GDRCOPY_LIBRARY REQUIRED
+      NAMES gdrapi
+    )
 
-# Set up required include directory flags, NVHPC CMake config only defined library directories
-string(REGEX REPLACE "(.+)(\/lib64)$" "\\1/include" NVHPC_CUDA_INCLUDE_DIR ${NVHPC_CUDA_LIBRARY_DIR})
-string(REGEX REPLACE "(.+)(\/lib64)$" "\\1/include" NVHPC_CUFFT_INCLUDE_DIR ${NVHPC_MATH_LIBRARY_DIR})
-string(REGEX REPLACE "(.+)(\/lib64)$" "\\1/include" NVHPC_CUTENSOR_INCLUDE_DIR ${NVHPC_MATH_LIBRARY_DIR})
+    message(STATUS "Found Cray GTL library: " ${CRAY_MPI_GTL_CUDA_LIBRARY})
+    message(STATUS "Found GDRCopy library: " ${GDRCOPY_LIBRARY})
+  endif()
 
-# Get NCCL library (with optional override)
-if (CUDECOMP_NCCL_HOME)
-  find_path(NCCL_INCLUDE_DIR REQUIRED
-    NAMES nccl.h
-    HINTS ${CUDECOMP_NCCL_HOME}/include
-  )
+  # HPC SDK
+  find_package(NVHPC REQUIRED COMPONENTS CUDA MATH)
 
-  find_library(NCCL_LIBRARY REQUIRED
-    NAMES nccl
-    HINTS ${CUDECOMP_NCCL_HOME}/lib
-  )
-else()
-  find_package(NVHPC REQUIRED COMPONENTS NCCL)
-  find_library(NCCL_LIBRARY
-    NAMES nccl
-    HINTS ${NVHPC_NCCL_LIBRARY_DIR}
-  )
-  string(REGEX REPLACE "(.+)(\/lib)$" "\\1/include" NCCL_INCLUDE_DIR ${NVHPC_NCCL_LIBRARY_DIR})
-endif()
+  # Set up required include directory flags, NVHPC CMake config only defined library directories
+  string(REGEX REPLACE "(.+)(\/lib64)$" "\\1/include" NVHPC_CUDA_INCLUDE_DIR ${NVHPC_CUDA_LIBRARY_DIR})
+  string(REGEX REPLACE "(.+)(\/lib64)$" "\\1/include" NVHPC_CUFFT_INCLUDE_DIR ${NVHPC_MATH_LIBRARY_DIR})
+  string(REGEX REPLACE "(.+)(\/lib64)$" "\\1/include" NVHPC_CUTENSOR_INCLUDE_DIR ${NVHPC_MATH_LIBRARY_DIR})
 
-message(STATUS "Using NCCL library: ${NCCL_LIBRARY}")
-
-if (CUDECOMP_ENABLE_NVSHMEM)
-  # Get NVSHMEM library (with optional override)
-  if (CUDECOMP_NVSHMEM_HOME)
-    find_path(NVSHMEM_INCLUDE_DIR REQUIRED
-      NAMES nvshmem.h
-      HINTS ${CUDECOMP_NVSHMEM_HOME}/include
+  # Get NCCL library (with optional override)
+  if (CUDECOMP_NCCL_HOME)
+    find_path(NCCL_INCLUDE_DIR REQUIRED
+      NAMES nccl.h
+      HINTS ${CUDECOMP_NCCL_HOME}/include
     )
 
-    find_path(NVSHMEM_LIBRARY_DIR REQUIRED
-      NAMES libnvshmem.a
-      HINTS ${CUDECOMP_NVSHMEM_HOME}/lib
+    find_library(NCCL_LIBRARY REQUIRED
+      NAMES nccl
+      HINTS ${CUDECOMP_NCCL_HOME}/lib
     )
   else()
-    find_package(NVHPC REQUIRED COMPONENTS NVSHMEM)
-    set(NVSHMEM_LIBRARY_DIR ${NVHPC_NVSHMEM_LIBRARY_DIR})
-    string(REGEX REPLACE "(.+)(\/lib)$" "\\1/include" NVSHMEM_INCLUDE_DIR ${NVHPC_NVSHMEM_LIBRARY_DIR})
+    find_package(NVHPC REQUIRED COMPONENTS NCCL)
+    find_library(NCCL_LIBRARY
+      NAMES nccl
+      HINTS ${NVHPC_NCCL_LIBRARY_DIR}
+    )
+    string(REGEX REPLACE "(.+)(\/lib)$" "\\1/include" NCCL_INCLUDE_DIR ${NVHPC_NCCL_LIBRARY_DIR})
   endif()
 
-  message(STATUS "Using NVSHMEM installation at: ${NVSHMEM_LIBRARY_DIR}")
+  message(STATUS "Using NCCL library: ${NCCL_LIBRARY}")
+
+  if (CUDECOMP_ENABLE_NVSHMEM)
+    # Get NVSHMEM library (with optional override)
+    if (CUDECOMP_NVSHMEM_HOME)
+      find_path(NVSHMEM_INCLUDE_DIR REQUIRED
+        NAMES nvshmem.h
+        HINTS ${CUDECOMP_NVSHMEM_HOME}/include
+      )
+
+      find_path(NVSHMEM_LIBRARY_DIR REQUIRED
+        NAMES libnvshmem.a
+        HINTS ${CUDECOMP_NVSHMEM_HOME}/lib
+      )
+    else()
+      find_package(NVHPC REQUIRED COMPONENTS NVSHMEM)
+      set(NVSHMEM_LIBRARY_DIR ${NVHPC_NVSHMEM_LIBRARY_DIR})
+      string(REGEX REPLACE "(.+)(\/lib)$" "\\1/include" NVSHMEM_INCLUDE_DIR ${NVHPC_NVSHMEM_LIBRARY_DIR})
+    endif()
+
+    message(STATUS "Using NVSHMEM installation at: ${NVSHMEM_LIBRARY_DIR}")
+
+  endif()
 
 endif()
 
@@ -174,128 +193,141 @@ set_target_properties(
              SOVERSION ${CUDECOMP_SOVERSION}
 )
 
-# Set NVCC flags for requested compute capability
-if (CMAKE_VERSION VERSION_LESS 3.18)
-  foreach(CUDA_CC ${CUDECOMP_CUDA_CC_LIST})
-    list(APPEND CUDA_CC_FLAGS -gencode=arch=compute_${CUDA_CC},code=sm_${CUDA_CC})
-  endforeach()
-  target_compile_options(cudecomp PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: ${CUDA_CC_FLAGS}>)
-else()
-  set_target_properties(cudecomp PROPERTIES CUDA_ARCHITECTURES "${CUDECOMP_CUDA_CC_LIST}")
-endif()
-
-target_compile_options(cudecomp PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-Wno-deprecated-gpu-targets >)
-
-target_sources(cudecomp
-  PRIVATE
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/autotune.cc
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/cuda_wrap.cc
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_kernels.cu
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_kernels_rdc.cu
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp.cc
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/graph.cc
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/nvml_wrap.cc
-  ${CMAKE_CURRENT_SOURCE_DIR}/src/performance.cc
-)
-
-set_source_files_properties(${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_kernels_rdc.cu PROPERTIES COMPILE_FLAGS -rdc=true)
-set_target_properties(cudecomp PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON)
-
-target_include_directories(cudecomp
-  PRIVATE
-  ${CMAKE_CURRENT_SOURCE_DIR}/include
-  ${MPI_CXX_INCLUDE_DIRS}
-  ${NVHPC_CUDA_INCLUDE_DIR}
-  ${NVHPC_CUTENSOR_INCLUDE_DIR}
-  ${NCCL_INCLUDE_DIR}
+target_sources(cudecomp PRIVATE src/cudecomp.cc)
+target_include_directories(cudecomp PUBLIC
+  $<BUILD_INTERFACE:${CMAKE_BINARY_DIR}/include>
+  $<INSTALL_INTERFACE:include>
 )
-
-target_link_libraries(cudecomp PUBLIC NVHPC::CUDART)
+target_include_directories(cudecomp PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
 target_link_libraries(cudecomp PUBLIC MPI::MPI_CXX)
-target_link_libraries(cudecomp PRIVATE NVHPC::CUTENSOR)
-target_link_libraries(cudecomp PRIVATE ${NCCL_LIBRARY})
-if (CRAY_CC_BIN)
-  # In Cray environments, add links to GTL and GDRCopy libs for CUDA-aware support
-  target_link_libraries(cudecomp PRIVATE ${CRAY_MPI_GTL_CUDA_LIBRARY})
-  target_link_libraries(cudecomp PRIVATE ${GDRCOPY_LIBRARY})
-endif()
+if (CUDECOMP_BUILD_CPU_ONLY)
+  target_sources(cudecomp PRIVATE src/cudecomp_kernels_cpu.cc)
+else()
+  # Set NVCC flags for requested compute capability
+  if (CMAKE_VERSION VERSION_LESS 3.18)
+    foreach(CUDA_CC ${CUDECOMP_CUDA_CC_LIST})
+      list(APPEND CUDA_CC_FLAGS -gencode=arch=compute_${CUDA_CC},code=sm_${CUDA_CC})
+    endforeach()
+    target_compile_options(cudecomp PRIVATE $<$<COMPILE_LANGUAGE:CUDA>: ${CUDA_CC_FLAGS}>)
+  else()
+    set_target_properties(cudecomp PROPERTIES CUDA_ARCHITECTURES "${CUDECOMP_CUDA_CC_LIST}")
+  endif()
 
-# Test for std::filesystem linking
-try_compile(
-  TEST_STD_FS_RESULT
-  ${CMAKE_BINARY_DIR}
-  ${CMAKE_CURRENT_SOURCE_DIR}/cmake/test_std_fs.cc
-)
-if (NOT TEST_STD_FS_RESULT)
-  message(STATUS "GCC toolchain version does not link std::filesystem symbols by default. Adding explicit -lstdc++fs link.")
-  target_link_libraries(cudecomp PRIVATE -lstdc++fs)
-endif()
+  target_compile_options(cudecomp PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-Wno-deprecated-gpu-targets >)
 
-if (CUDECOMP_ENABLE_NVTX)
-  target_compile_definitions(cudecomp PRIVATE ENABLE_NVTX)
-endif()
+  target_sources(cudecomp
+    PRIVATE
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/autotune.cc
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/cuda_wrap.cc
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_kernels.cu
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_kernels_rdc.cu
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/graph.cc
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/nvml_wrap.cc
+    ${CMAKE_CURRENT_SOURCE_DIR}/src/performance.cc
+  )
+
+  set_source_files_properties(${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_kernels_rdc.cu PROPERTIES COMPILE_FLAGS -rdc=true)
+  set_target_properties(cudecomp PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 
-if (CUDECOMP_ENABLE_NVSHMEM)
-  target_compile_definitions(cudecomp PRIVATE ENABLE_NVSHMEM)
   target_include_directories(cudecomp
     PRIVATE
-    ${NVSHMEM_INCLUDE_DIR}
+    ${CMAKE_CURRENT_SOURCE_DIR}/include
+    ${MPI_CXX_INCLUDE_DIRS}
+    ${NVHPC_CUDA_INCLUDE_DIR}
+    ${NVHPC_CUTENSOR_INCLUDE_DIR}
+    ${NCCL_INCLUDE_DIR}
   )
 
-  # Get NVSHMEM version from header
-  if (EXISTS ${NVSHMEM_INCLUDE_DIR}/non_abi/nvshmem_version.h)
-    file(READ ${NVSHMEM_INCLUDE_DIR}/non_abi/nvshmem_version.h NVSHMEM_VERSION_RAW)
-  elseif (EXISTS ${NVSHMEM_INCLUDE_DIR}/nvshmem_version.h)
-    file(READ ${NVSHMEM_INCLUDE_DIR}/nvshmem_version.h NVSHMEM_VERSION_RAW)
-  else()
-    file(READ ${NVSHMEM_INCLUDE_DIR}/common/nvshmem_version.h NVSHMEM_VERSION_RAW)
+  target_link_libraries(cudecomp PUBLIC NVHPC::CUDART)
+  target_link_libraries(cudecomp PRIVATE NVHPC::CUTENSOR)
+  target_link_libraries(cudecomp PRIVATE ${NCCL_LIBRARY})
+  if (CRAY_CC_BIN)
+    # In Cray environments, add links to GTL and GDRCopy libs for CUDA-aware support
+    target_link_libraries(cudecomp PRIVATE ${CRAY_MPI_GTL_CUDA_LIBRARY})
+    target_link_libraries(cudecomp PRIVATE ${GDRCOPY_LIBRARY})
   endif()
-  string(REGEX MATCH "NVSHMEM_VENDOR_MAJOR_VERSION ([0-9]*)" _ ${NVSHMEM_VERSION_RAW})
-  list(APPEND NVSHMEM_VERSION ${CMAKE_MATCH_1})
-  string(REGEX MATCH "NVSHMEM_VENDOR_MINOR_VERSION ([0-9]*)" _ ${NVSHMEM_VERSION_RAW})
-  list(APPEND NVSHMEM_VERSION ${CMAKE_MATCH_1})
-  list(JOIN NVSHMEM_VERSION "." NVSHMEM_VERSION)
-
-  if (NVSHMEM_VERSION VERSION_LESS "2.6")
-    message(FATAL_ERROR "NVSHMEM versions earlier than 2.6.0 are not supported by cuDecomp.")
-  else()
-    target_link_libraries(cudecomp PRIVATE ${NVSHMEM_LIBRARY_DIR}/libnvshmem_host.so)
-    target_link_libraries(cudecomp PRIVATE ${NVSHMEM_LIBRARY_DIR}/libnvshmem_device.a)
+
+  # Test for std::filesystem linking
+  try_compile(
+    TEST_STD_FS_RESULT
+    ${CMAKE_BINARY_DIR}
+    ${CMAKE_CURRENT_SOURCE_DIR}/cmake/test_std_fs.cc
+  )
+  if (NOT TEST_STD_FS_RESULT)
+    message(STATUS "GCC toolchain version does not link std::filesystem symbols by default. Adding explicit -lstdc++fs link.")
+    target_link_libraries(cudecomp PRIVATE -lstdc++fs)
   endif()
 
-  if (NVSHMEM_VERSION VERSION_LESS "2.7")
-    # NVSHMEM versions before 2.7 will export NCCL symbols erroneously, need to define this flag
-    target_compile_definitions(cudecomp PRIVATE NVSHMEM_USE_NCCL)
-    target_link_libraries(cudecomp PUBLIC -L${NVHPC_CUDA_LIBRARY_DIR}/stubs -lcuda)
+  if (CUDECOMP_ENABLE_NVTX)
+    target_compile_definitions(cudecomp PRIVATE ENABLE_NVTX)
   endif()
 
-  if (NVSHMEM_VERSION VERSION_LESS "2.11")
-    target_link_libraries(cudecomp PUBLIC -L${NVHPC_CUDA_LIBRARY_DIR}/stubs -lnvidia-ml)
+  if (CUDECOMP_ENABLE_NVSHMEM)
+    target_compile_definitions(cudecomp PRIVATE ENABLE_NVSHMEM)
+    target_include_directories(cudecomp
+      PRIVATE
+      ${NVSHMEM_INCLUDE_DIR}
+    )
+
+    # Get NVSHMEM version from header
+    if (EXISTS ${NVSHMEM_INCLUDE_DIR}/non_abi/nvshmem_version.h)
+      file(READ ${NVSHMEM_INCLUDE_DIR}/non_abi/nvshmem_version.h NVSHMEM_VERSION_RAW)
+    elseif (EXISTS ${NVSHMEM_INCLUDE_DIR}/nvshmem_version.h)
+      file(READ ${NVSHMEM_INCLUDE_DIR}/nvshmem_version.h NVSHMEM_VERSION_RAW)
+    else()
+      file(READ ${NVSHMEM_INCLUDE_DIR}/common/nvshmem_version.h NVSHMEM_VERSION_RAW)
+    endif()
+    string(REGEX MATCH "NVSHMEM_VENDOR_MAJOR_VERSION ([0-9]*)" _ ${NVSHMEM_VERSION_RAW})
+    list(APPEND NVSHMEM_VERSION ${CMAKE_MATCH_1})
+    string(REGEX MATCH "NVSHMEM_VENDOR_MINOR_VERSION ([0-9]*)" _ ${NVSHMEM_VERSION_RAW})
+    list(APPEND NVSHMEM_VERSION ${CMAKE_MATCH_1})
+    list(JOIN NVSHMEM_VERSION "." NVSHMEM_VERSION)
+
+    if (NVSHMEM_VERSION VERSION_LESS "2.6")
+      message(FATAL_ERROR "NVSHMEM versions earlier than 2.6.0 are not supported by cuDecomp.")
+    else()
+      target_link_libraries(cudecomp PRIVATE ${NVSHMEM_LIBRARY_DIR}/libnvshmem_host.so)
+      target_link_libraries(cudecomp PRIVATE ${NVSHMEM_LIBRARY_DIR}/libnvshmem_device.a)
+    endif()
+
+    if (NVSHMEM_VERSION VERSION_LESS "2.7")
+      # NVSHMEM versions before 2.7 will export NCCL symbols erroneously, need to define this flag
+      target_compile_definitions(cudecomp PRIVATE NVSHMEM_USE_NCCL)
+      target_link_libraries(cudecomp PUBLIC -L${NVHPC_CUDA_LIBRARY_DIR}/stubs -lcuda)
+    endif()
+
+    if (NVSHMEM_VERSION VERSION_LESS "2.11")
+      target_link_libraries(cudecomp PUBLIC -L${NVHPC_CUDA_LIBRARY_DIR}/stubs -lnvidia-ml)
+    endif()
   endif()
+
 endif()
 
+configure_file(include/cudecomp_config.h.in ${CMAKE_BINARY_DIR}/include/cudecomp_config.h)
 set_target_properties(
   cudecomp
   PROPERTIES PUBLIC_HEADER
-             "${CMAKE_CURRENT_SOURCE_DIR}/include/cudecomp.h;${CMAKE_CURRENT_SOURCE_DIR}/include/cudecomp_version.h"
+             "${CMAKE_CURRENT_SOURCE_DIR}/include/cudecomp.h;${CMAKE_CURRENT_SOURCE_DIR}/include/cudecomp_version.h;${CMAKE_BINARY_DIR}/include/cudecomp_config.h"
 )
 configure_file(${CMAKE_CURRENT_SOURCE_DIR}/include/cudecomp.h ${CMAKE_BINARY_DIR}/include/cudecomp.h)
 configure_file(${CMAKE_CURRENT_SOURCE_DIR}/include/cudecomp_version.h ${CMAKE_BINARY_DIR}/include/cudecomp_version.h)
 
 install(
   TARGETS cudecomp
-  LIBRARY DESTINATION ${CMAKE_INSTALL_PREFIX}/lib
-  PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_PREFIX}/include
+  LIBRARY DESTINATION lib
+  PUBLIC_HEADER DESTINATION include
 )
 
 # Building Fortran shared lib and module
 if (CUDECOMP_BUILD_FORTRAN)
-  # Creating -gpu argument string for Fortran files
-  foreach(CUDA_CC ${CUDECOMP_CUDA_CC_LIST})
-    list(APPEND CUF_GPU_ARG "cc${CUDA_CC}")
-  endforeach()
-  list(APPEND CUF_GPU_ARG "cuda${NVHPC_CUDA_VERSION}")
-  list(JOIN CUF_GPU_ARG "," CUF_GPU_ARG)
+  if (NOT CUDECOMP_BUILD_CPU_ONLY)
+    # Creating -gpu argument string for Fortran files
+    foreach(CUDA_CC ${CUDECOMP_CUDA_CC_LIST})
+      list(APPEND CUF_GPU_ARG "cc${CUDA_CC}")
+    endforeach()
+    list(APPEND CUF_GPU_ARG "cuda${NVHPC_CUDA_VERSION}")
+    list(JOIN CUF_GPU_ARG "," CUF_GPU_ARG)
+  endif()
 
   add_library(cudecomp_fort SHARED)
   set_target_properties(
@@ -306,34 +338,41 @@ if (CUDECOMP_BUILD_FORTRAN)
                VERSION ${PROJECT_VERSION}
                SOVERSION ${CUDECOMP_SOVERSION}
   )
-  target_compile_options(cudecomp_fort PRIVATE $<$<COMPILE_LANGUAGE:Fortran>:-cpp -cuda -gpu=${CUF_GPU_ARG}>)
-  target_sources(
-    cudecomp_fort
-    PRIVATE
-    ${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_m.cuf
-  )
-  set_source_files_properties(${CMAKE_CURRENT_SOURCE_DIR}/src/cudecomp_m.cuf PROPERTIES LANGUAGE Fortran)
+  if (CUDECOMP_BUILD_CPU_ONLY)
+    # A standard extension lets GNU Fortran preprocess the shared binding source.
+    configure_file(src/cudecomp_m.cuf ${CMAKE_CURRENT_BINARY_DIR}/cudecomp_m.F90 COPYONLY)
+    target_sources(cudecomp_fort PRIVATE ${CMAKE_CURRENT_BINARY_DIR}/cudecomp_m.F90)
+    target_compile_definitions(cudecomp_fort PRIVATE CUDECOMP_BUILD_CPU_ONLY=1)
+    target_compile_options(cudecomp_fort PRIVATE $<$<Fortran_COMPILER_ID:GNU>:-ffree-line-length-none>)
+  else()
+    target_compile_options(cudecomp_fort PRIVATE $<$<COMPILE_LANGUAGE:Fortran>:-cpp -cuda -gpu=${CUF_GPU_ARG}>)
+    target_sources(cudecomp_fort PRIVATE src/cudecomp_m.cuf)
+    set_source_files_properties(src/cudecomp_m.cuf PROPERTIES LANGUAGE Fortran)
+  endif()
 
   target_link_libraries(cudecomp_fort PUBLIC MPI::MPI_Fortran)
 
   install(
     TARGETS cudecomp_fort
-    LIBRARY DESTINATION ${CMAKE_INSTALL_PREFIX}/lib
+    LIBRARY DESTINATION lib
    )
    # Install cuDecomp modules
    install(
      FILES ${CMAKE_BINARY_DIR}/include/cudecomp.mod
            ${CMAKE_BINARY_DIR}/include/cudecomp_internal.mod
-     DESTINATION ${CMAKE_INSTALL_PREFIX}/include
+     DESTINATION include
    )
 endif()
 
 if (CUDECOMP_BUILD_TESTS OR CUDECOMP_BUILD_EXTRAS)
   enable_testing()
-  add_subdirectory(tests/cc)
+  if (NOT CUDECOMP_BUILD_CPU_ONLY)
+    add_subdirectory(tests/cc)
+
+    if (CUDECOMP_BUILD_FORTRAN)
+      add_subdirectory(tests/fortran)
+    endif()
 
-  if (CUDECOMP_BUILD_FORTRAN)
-    add_subdirectory(tests/fortran)
   endif()
 
   if (CUDECOMP_BUILD_TESTS)
@@ -341,10 +380,9 @@ if (CUDECOMP_BUILD_TESTS OR CUDECOMP_BUILD_EXTRAS)
   endif()
 endif()
 
-if (CUDECOMP_BUILD_EXTRAS)
-  add_subdirectory(benchmark)
-
+if (CUDECOMP_BUILD_EXTRAS AND NOT CUDECOMP_BUILD_CPU_ONLY)
   add_subdirectory(examples/cc/basic_usage)
+  add_subdirectory(benchmark)
   add_subdirectory(examples/cc/taylor_green)
 
   if (CUDECOMP_BUILD_FORTRAN)
diff --git a/README.md b/README.md
index 8263b0f..f4dcc90 100644
--- a/README.md
+++ b/README.md
@@ -27,8 +27,8 @@ $ cmake -DCUDECOMP_BUILD_EXTRAS=1 -DCUDECOMP_ENABLE_NVSHMEM=1 ..
 ```
 
 
-### Dependencies
-We strongly recommend building this library using NVHPC SDK compilers and libraries, as the SDK contains all required dependencies for this library and is the focus of our testing. Fortran features are only supported using NVHPC SDK compilers.
+### GPU dependencies
+We strongly recommend building this library using NVHPC SDK compilers and libraries, as the SDK contains all required dependencies for this library and is the focus of our testing. GPU Fortran bindings require NVHPC SDK compilers.
 
 One exception is cuDecomp builds using NVSHMEM versions older than v3.0, which require the use of a bootstrapping layer that depends on your MPI distribution. The NVSHMEM library packaged within NVHPC SDK
 supports OpenMPI only. If you require usage of a different MPI implementation (e.g. Spectrum MPI or Cray MPICH), you need to either build
@@ -39,4 +39,3 @@ Additionally, this library utilizes CUDA-aware MPI and is only compatible with M
 
 ## License
 This library is released under an Apache 2.0 license, which can be found in [LICENSE](LICENSE).
-
diff --git a/docs/cpu_only.rst b/docs/cpu_only.rst
new file mode 100644
index 0000000..1ef7353
--- /dev/null
+++ b/docs/cpu_only.rst
@@ -0,0 +1,18 @@
+CPU-only build
+==============
+
+Configure with ``-DCUDECOMP_BUILD_CPU_ONLY=ON`` for a limited, toolkit-free
+debug/test build requiring C++17 and MPI (plus a Fortran compiler for the optional
+bindings). This mode is intended for debugging and testing only, not for production
+use. The API is unchanged except for the following differences:
+
+* Library, function, and Fortran module names are unchanged. Use separate CPU/GPU
+  build directories and installation prefixes, and only one variant per process.
+* Buffers and ``cudecompMalloc`` allocations reside in host memory. Operations
+  complete synchronously; stream arguments are ignored.
+* Only MPI backends are supported. Set both ``pdims`` explicitly, with their product
+  equal to the communicator size. Autotuning is unsupported; pass null autotuning
+  options or disable tuning.
+* Fortran bindings use host arrays and standard C pointers instead of device
+  arrays/pointers.
+* CUDA Graphs, performance reports, NCCL, NVSHMEM, and NVTX are unsupported.
diff --git a/docs/index.rst b/docs/index.rst
index 12a1fbf..e4e9d7a 100644
--- a/docs/index.rst
+++ b/docs/index.rst
@@ -27,9 +27,10 @@ Table of Contents
 
    overview
    basic_usage
-   cuda_graph_capture
    autotuning
+   cuda_graph_capture
    nvshmem
+   cpu_only
    api
    env_vars
 
diff --git a/include/cudecomp.h b/include/cudecomp.h
index 00788b3..c5a411f 100644
--- a/include/cudecomp.h
+++ b/include/cudecomp.h
@@ -27,7 +27,13 @@
 #include <stdbool.h>
 #include <stdint.h>
 
+#include "cudecomp_config.h"
+#if CUDECOMP_BUILD_CPU_ONLY
+// Keep stream signatures compatible without requiring CUDA headers.
+typedef struct CUstream_st* cudaStream_t;
+#else
 #include <cuda_runtime.h>
+#endif
 #include <mpi.h>
 
 #include "cudecomp_version.h"
diff --git a/include/cudecomp_config.h.in b/include/cudecomp_config.h.in
new file mode 100644
index 0000000..67f1c71
--- /dev/null
+++ b/include/cudecomp_config.h.in
@@ -0,0 +1,10 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#ifndef CUDECOMP_CONFIG_H
+#define CUDECOMP_CONFIG_H
+
+#cmakedefine01 CUDECOMP_BUILD_CPU_ONLY
+
+#endif
diff --git a/include/internal/checks.h b/include/internal/checks.h
index 67a3fc9..683f38c 100644
--- a/include/internal/checks.h
+++ b/include/internal/checks.h
@@ -24,13 +24,18 @@
 #include <sstream>
 #include <string>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda.h>
 #include <cuda_runtime.h>
 #include <cufft.h>
 #include <cutensor.h>
+#endif
 #include <mpi.h>
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <nccl.h>
 #include <nvml.h>
+#endif
 
 #include "internal/exceptions.h"
 
diff --git a/include/internal/comm_routines.h b/include/internal/comm_routines.h
index a347238..105a2d5 100644
--- a/include/internal/comm_routines.h
+++ b/include/internal/comm_routines.h
@@ -27,10 +27,15 @@
 #include <string>
 #include <vector>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda/std/complex>
 #include <cuda_runtime.h>
+#endif
 #include <mpi.h>
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <nccl.h>
+#endif
 #ifdef ENABLE_NVSHMEM
 #include <nvshmem.h>
 #include <nvshmemx.h>
@@ -39,6 +44,7 @@
 #include "internal/checks.h"
 #include "internal/cudecomp_kernels.h"
 #include "internal/nvtx.h"
+#include "internal/utils.h"
 
 namespace cudecomp {
 
@@ -48,8 +54,8 @@ inline constexpr uint64_t nvshmem_signal_set_value = 1; // Graph-stable source f
 
 static inline MPI_Datatype getMpiDataType(float) { return MPI_FLOAT; }
 static inline MPI_Datatype getMpiDataType(double) { return MPI_DOUBLE; }
-static inline MPI_Datatype getMpiDataType(cuda::std::complex<float>) { return MPI_C_FLOAT_COMPLEX; }
-static inline MPI_Datatype getMpiDataType(cuda::std::complex<double>) { return MPI_C_DOUBLE_COMPLEX; }
+static inline MPI_Datatype getMpiDataType(cudecomp::complex<float>) { return MPI_C_FLOAT_COMPLEX; }
+static inline MPI_Datatype getMpiDataType(cudecomp::complex<double>) { return MPI_C_DOUBLE_COMPLEX; }
 template <typename T> static inline MPI_Datatype getMpiDataType() { return getMpiDataType(T(0)); }
 
 static inline bool canUseMpiAlltoall(const std::vector<comm_count_t>& send_counts,
@@ -271,9 +277,11 @@ static void cudecompAlltoall(const cudecompHandle_t& handle, const cudecompGridD
                              cudaStream_t stream, cudecompTransposePerformanceSample* current_sample = nullptr) {
   nvtx::rangePush("cudecompAlltoall");
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable) {
     CHECK_CUDA(cudaEventRecord(current_sample->alltoall_start_events[current_sample->alltoall_timing_count], stream));
   }
+#endif
 
 #ifdef ENABLE_NVSHMEM
   if (transposeBackendRequiresMpi(grid_desc->config.transpose_comm_backend)) {
@@ -298,6 +306,7 @@ static void cudecompAlltoall(const cudecompHandle_t& handle, const cudecompGridD
     THROW_NOT_SUPPORTED("build does not support NVSHMEM communication backends.");
 #endif
   }
+#if !CUDECOMP_BUILD_CPU_ONLY
   case CUDECOMP_TRANSPOSE_COMM_NCCL: {
     auto& comm_info = (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info : grid_desc->col_comm_info;
     // For fully intra-group alltoall, use distinct NCCL local comm instead of global comm as it is faster.
@@ -326,17 +335,17 @@ static void cudecompAlltoall(const cudecompHandle_t& handle, const cudecompGridD
     CHECK_NCCL(ncclGroupEnd());
     break;
   }
+#endif
   case CUDECOMP_TRANSPOSE_COMM_MPI_P2P: {
     std::vector<MPI_Request> reqs(2 * send_counts.size(), MPI_REQUEST_NULL);
-    CHECK_CUDA(cudaStreamSynchronize(stream));
+    streamSynchronize(stream);
 
     auto comm =
         (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info.mpi_comm : grid_desc->col_comm_info.mpi_comm;
     int self_rank = (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info.rank : grid_desc->col_comm_info.rank;
 
-    // Self-copy with cudaMemcpy
-    CHECK_CUDA(cudaMemcpyAsync(recv_buff + recv_offsets[self_rank], send_buff + send_offsets[self_rank],
-                               send_counts[self_rank] * sizeof(T), cudaMemcpyDeviceToDevice, stream));
+    copyBuffer(recv_buff + recv_offsets[self_rank], send_buff + send_offsets[self_rank],
+               send_counts[self_rank] * sizeof(T), stream);
 
     for (int i = 1; i < recv_counts.size(); ++i) {
       int src_rank, dst_rank;
@@ -365,7 +374,7 @@ static void cudecompAlltoall(const cudecompHandle_t& handle, const cudecompGridD
     break;
   }
   case CUDECOMP_TRANSPOSE_COMM_MPI_A2A: {
-    CHECK_CUDA(cudaStreamSynchronize(stream));
+    streamSynchronize(stream);
     auto comm =
         (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info.mpi_comm : grid_desc->col_comm_info.mpi_comm;
     int self_rank = (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info.rank : grid_desc->col_comm_info.rank;
@@ -379,9 +388,8 @@ static void cudecompAlltoall(const cudecompHandle_t& handle, const cudecompGridD
       int32_t rc = static_cast<int32_t>(recv_counts[0]);
       CHECK_MPI(MPI_Alltoall(send_buff, sc, getMpiDataType<T>(), recv_buff, rc, getMpiDataType<T>(), comm));
     } else {
-      // Self-copy with cudaMemcpy
-      CHECK_CUDA(cudaMemcpyAsync(recv_buff + recv_offsets[self_rank], send_buff + send_offsets[self_rank],
-                                 send_counts[self_rank] * sizeof(T), cudaMemcpyDeviceToDevice, stream));
+      copyBuffer(recv_buff + recv_offsets[self_rank], send_buff + send_offsets[self_rank],
+                 send_counts[self_rank] * sizeof(T), stream);
 
       // Convert count/offset args to int32
       std::vector<int32_t> send_counts_i32(send_counts.size());
@@ -420,10 +428,12 @@ static void cudecompAlltoall(const cudecompHandle_t& handle, const cudecompGridD
   }
   }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable) {
     CHECK_CUDA(cudaEventRecord(current_sample->alltoall_end_events[current_sample->alltoall_timing_count], stream));
     current_sample->alltoall_timing_count++;
   }
+#endif
 
   nvtx::rangePop();
 }
@@ -450,6 +460,7 @@ cudecompAlltoallPipelined(const cudecompHandle_t& handle, const cudecompGridDesc
   }
   nvtx::rangePush(os.str());
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   int self_rank = (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info.rank : grid_desc->col_comm_info.rank;
   if (handle->performance_report_enable && src_ranks[0] != self_rank) {
     // Note: skipping self-copy for timing as it should be overlapped
@@ -457,6 +468,7 @@ cudecompAlltoallPipelined(const cudecompHandle_t& handle, const cudecompGridDesc
     CHECK_CUDA(cudaEventRecord(current_sample->alltoall_start_events[current_sample->alltoall_timing_count],
                                handle->streams[0]));
   }
+#endif
 
 #ifdef ENABLE_NVSHMEM
   if (transposeBackendRequiresMpi(grid_desc->config.transpose_comm_backend)) {
@@ -537,6 +549,7 @@ cudecompAlltoallPipelined(const cudecompHandle_t& handle, const cudecompGridDesc
     THROW_NOT_SUPPORTED("build does not support NVSHMEM communication backends.");
 #endif
   }
+#if !CUDECOMP_BUILD_CPU_ONLY
   case CUDECOMP_TRANSPOSE_COMM_NCCL_PL: {
     const auto& comm_info = (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info : grid_desc->col_comm_info;
     // For fully intra-group alltoall, use distinct NCCL local comm instead of global comm as it is faster.
@@ -589,6 +602,7 @@ cudecompAlltoallPipelined(const cudecompHandle_t& handle, const cudecompGridDesc
     }
     break;
   }
+#endif
   case CUDECOMP_TRANSPOSE_COMM_MPI_P2P_PL: {
     auto comm =
         (comm_axis == CUDECOMP_COMM_ROW) ? grid_desc->row_comm_info.mpi_comm : grid_desc->col_comm_info.mpi_comm;
@@ -599,11 +613,12 @@ cudecompAlltoallPipelined(const cudecompHandle_t& handle, const cudecompGridDesc
       int src_rank = src_ranks[i];
       int dst_rank = dst_ranks[i];
       if (src_rank == self_rank) {
-        // Self-copy with cudaMemcpy
-        CHECK_CUDA(cudaMemcpyAsync(recv_buff + recv_offsets[self_rank], send_buff + send_offsets[self_rank],
-                                   send_counts[self_rank] * sizeof(T), cudaMemcpyDeviceToDevice, stream));
+        copyBuffer(recv_buff + recv_offsets[self_rank], send_buff + send_offsets[self_rank],
+                   send_counts[self_rank] * sizeof(T), stream);
       } else {
+#if !CUDECOMP_BUILD_CPU_ONLY
         CHECK_CUDA(cudaEventSynchronize(grid_desc->events[dst_rank]));
+#endif
 
         if (send_counts[dst_rank] != 0) {
           checkMpiInt32Limit(send_counts[dst_rank], grid_desc->config.transpose_comm_backend);
@@ -629,11 +644,13 @@ cudecompAlltoallPipelined(const cudecompHandle_t& handle, const cudecompGridDesc
   }
   }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable && src_ranks[0] != self_rank) {
     CHECK_CUDA(cudaEventRecord(current_sample->alltoall_end_events[current_sample->alltoall_timing_count],
                                handle->streams[0]));
     current_sample->alltoall_timing_count++;
   }
+#endif
   nvtx::rangePop();
 }
 
@@ -647,9 +664,11 @@ static void cudecompSendRecvPair(const cudecompHandle_t& handle, const cudecompG
                                  cudecompHaloPerformanceSample* current_sample = nullptr) {
   nvtx::rangePush("cudecompSendRecvPair");
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable && current_sample) {
     CHECK_CUDA(cudaEventRecord(current_sample->sendrecv_start_event, stream));
   }
+#endif
 
 #ifdef ENABLE_NVSHMEM
   if (haloBackendRequiresMpi(grid_desc->config.halo_comm_backend)) {
@@ -690,6 +709,7 @@ static void cudecompSendRecvPair(const cudecompHandle_t& handle, const cudecompG
     THROW_NOT_SUPPORTED("build does not support NVSHMEM communication backends.");
 #endif
   }
+#if !CUDECOMP_BUILD_CPU_ONLY
   case CUDECOMP_HALO_COMM_NCCL: {
     auto comm = *grid_desc->nccl_comm;
     CHECK_NCCL(ncclGroupStart());
@@ -712,15 +732,14 @@ static void cudecompSendRecvPair(const cudecompHandle_t& handle, const cudecompG
     CHECK_NCCL(ncclGroupEnd());
     break;
   }
+#endif
   case CUDECOMP_HALO_COMM_MPI: {
     auto comm = handle->mpi_comm;
     std::vector<MPI_Request> reqs(2 * send_counts.size(), MPI_REQUEST_NULL);
-    CHECK_CUDA(cudaStreamSynchronize(stream));
+    streamSynchronize(stream);
     for (int i = 0; i < send_counts.size(); ++i) {
       if (peer_ranks[i] == handle->rank) {
-        // Self-copy with cudaMemcpy
-        CHECK_CUDA(cudaMemcpyAsync(recv_buff + recv_offsets[i], send_buff + send_offsets[i], send_counts[i] * sizeof(T),
-                                   cudaMemcpyDeviceToDevice, stream));
+        copyBuffer(recv_buff + recv_offsets[i], send_buff + send_offsets[i], send_counts[i] * sizeof(T), stream);
       } else {
         if (recv_counts[(i + 1) % 2] != 0 && peer_ranks[(i + 1) % 2] != -1) {
           checkMpiInt32Limit(recv_counts[(i + 1) % 2], grid_desc->config.halo_comm_backend);
@@ -742,12 +761,10 @@ static void cudecompSendRecvPair(const cudecompHandle_t& handle, const cudecompG
   }
   case CUDECOMP_HALO_COMM_MPI_BLOCKING: {
     auto comm = handle->mpi_comm;
-    CHECK_CUDA(cudaStreamSynchronize(stream));
+    streamSynchronize(stream);
     for (int i = 0; i < send_counts.size(); ++i) {
       if (peer_ranks[i] == handle->rank) {
-        // Self-copy with cudaMemcpy
-        CHECK_CUDA(cudaMemcpyAsync(recv_buff + recv_offsets[i], send_buff + send_offsets[i], send_counts[i] * sizeof(T),
-                                   cudaMemcpyDeviceToDevice, stream));
+        copyBuffer(recv_buff + recv_offsets[i], send_buff + send_offsets[i], send_counts[i] * sizeof(T), stream);
       } else {
         MPI_Request r = MPI_REQUEST_NULL;
         if (recv_counts[(i + 1) % 2] != 0 && peer_ranks[(i + 1) % 2] != -1) {
@@ -772,9 +789,11 @@ static void cudecompSendRecvPair(const cudecompHandle_t& handle, const cudecompG
   }
   }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable && current_sample) {
     CHECK_CUDA(cudaEventRecord(current_sample->sendrecv_end_event, stream));
   }
+#endif
 
   nvtx::rangePop();
 }
diff --git a/include/internal/common.h b/include/internal/common.h
index b886352..f887491 100644
--- a/include/internal/common.h
+++ b/include/internal/common.h
@@ -23,6 +23,7 @@
 #include <complex>
 #include <cstddef>
 #include <cstdint>
+#include <cstring>
 #include <functional>
 #include <map>
 #include <memory>
@@ -35,11 +36,16 @@
 #include <utility>
 #include <vector>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda_runtime.h>
 #include <cutensor.h>
+#endif
 #include <mpi.h>
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <nccl.h>
 #include <nvml.h>
+#endif
 #ifdef ENABLE_NVSHMEM
 #include <nvshmem.h>
 #include <nvshmemx.h>
@@ -47,8 +53,11 @@
 
 #include "cudecomp.h"
 #include "internal/checks.h"
+#include "internal/utils.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include "internal/graph.h"
 #include "internal/raii_wrappers.h"
+#endif
 
 namespace cudecomp {
 #if NVML_API_VERSION >= 12 && CUDART_VERSION >= 12040
@@ -56,7 +65,9 @@ typedef std::pair<std::array<unsigned char, NVML_GPU_FABRIC_UUID_LEN>, unsigned
 #else
 typedef std::pair<std::array<unsigned char, 1>, unsigned int> mnnvl_info;
 #endif
+#if !CUDECOMP_BUILD_CPU_ONLY
 typedef std::shared_ptr<ncclComm_t> ncclComm;
+#endif
 struct nvshmemRuntimeState {
 #ifdef ENABLE_NVSHMEM
   ~nvshmemRuntimeState() noexcept { finalize(); }
@@ -114,15 +125,18 @@ struct cudecompHandle {
   int32_t clique_rank;                      // MPI rank
   int32_t clique_nranks;                    // MPI size
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   // Entries for NCCL management
   cudecomp::ncclComm nccl_comm;       // NCCL communicator (global)
   cudecomp::ncclComm nccl_local_comm; // NCCL communicator (intra-node, or intra-clique on MNNVL systems)
-  bool nccl_enable_ubr = false;       // Flag to control NCCL user buffer registration usage
   std::unordered_map<void*, std::vector<std::pair<cudecomp::ncclComm, void*>>>
       nccl_ubr_handles; // map of allocated buffer address to NCCL registration handle(s)
 
   std::vector<cudecomp::cudaStream> streams; // internal streams for concurrent scheduling
 
+#endif
+  bool nccl_enable_ubr = false; // Always false in CPU-only builds.
+
   // Automatic workspaces are split by allocation domain but share one execution stream so operations submitted with
   // a null workspace cannot race with each other.
   struct ManagedWorkspace {
@@ -135,11 +149,14 @@ struct cudecompHandle {
   };
   ManagedWorkspace ordinary_workspace;
   ManagedWorkspace nvshmem_workspace;
+#if !CUDECOMP_BUILD_CPU_ONLY
   std::unique_ptr<cudecomp::cudaStream> workspace_stream;
   std::unique_ptr<cudecomp::cudaEvent> workspace_ingress_event;
   std::unique_ptr<cudecomp::cudaEvent> workspace_egress_event;
+#endif
   std::mutex workspace_mutex;
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 #if CUTENSOR_MAJOR >= 2
   cutensorHandle_t cutensor_handle = nullptr;            // cuTENSOR handle;
   cutensorPlanPreference_t cutensor_plan_pref = nullptr; // cuTENSOR plan preference;
@@ -148,6 +165,7 @@ struct cudecompHandle {
   cutensorHandle_t cutensor_handle; // cuTENSOR handle;
 #endif
 
+#endif
   std::vector<std::array<char, MPI_MAX_PROCESSOR_NAME>> hostnames; // list of hostnames by rank
   std::vector<int32_t> rank_to_local_rank;                         // list of local rank mappings
 
@@ -227,6 +245,10 @@ struct cudecompCommInfo {
   bool mnnvl_active = false; // flag to indicate whether communicator has MNNVL connections
 };
 
+// CPU communication only passes null sample pointers; timing resources are GPU-only.
+struct cudecompTransposePerformanceSample;
+struct cudecompHaloPerformanceSample;
+#if !CUDECOMP_BUILD_CPU_ONLY
 // Structure to contain data for transpose performance sample
 struct cudecompTransposePerformanceSample {
   cudecomp::cudaEventTimed transpose_start_event;
@@ -262,6 +284,8 @@ struct cudecompHaloPerformanceSampleCollection {
   int32_t warmup_count = 0;
 };
 
+#endif
+
 // cuDecomp grid descriptor containing grid-specific information
 struct cudecompGridDesc {
   ~cudecompGridDesc() noexcept {
@@ -282,14 +306,17 @@ struct cudecompGridDesc {
   cudecompCommInfo row_comm_info; // row communicator information
   cudecompCommInfo col_comm_info; // column communicator information
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   std::vector<cudecomp::cudaEvent> events; // CUDA events used for scheduling
   cudecomp::cudaEvent nvshmem_sync_event;  // NVSHMEM event used for synchronization
 
+#endif
 #ifdef ENABLE_NVSHMEM
   int* nvshmem_block_counters = nullptr;    // device memory counters for SM alltoallv last-block detection
   cudecomp::nvshmemRuntime nvshmem_runtime; // Shared reference to initialized NVSHMEM runtime
 #endif
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   cudecomp::graphCache graph_cache; // CUDA graph cache
 
   cudecomp::ncclComm nccl_comm; // NCCL communicator (global), shared from handle
@@ -313,11 +340,27 @@ struct cudecompGridDesc {
                      cudecompHaloPerformanceSampleCollection>
       halo_perf_samples_map;
 
+#endif
   bool initialized = false;
 };
 
 namespace cudecomp {
 
+// CPU local work is synchronous; GPU MPI paths must retain their stream ordering.
+static inline void streamSynchronize(cudaStream_t stream) {
+#if !CUDECOMP_BUILD_CPU_ONLY
+  CHECK_CUDA(cudaStreamSynchronize(stream));
+#endif
+}
+
+static inline void copyBuffer(void* dst, const void* src, size_t bytes, cudaStream_t stream) {
+#if CUDECOMP_BUILD_CPU_ONLY
+  if (bytes) { std::memcpy(dst, src, bytes); }
+#else
+  CHECK_CUDA(cudaMemcpyAsync(dst, src, bytes, cudaMemcpyDeviceToDevice, stream));
+#endif
+}
+
 using comm_count_t = int64_t;
 
 enum cudecompCommAxis { CUDECOMP_COMM_COL = 0, CUDECOMP_COMM_ROW = 1 };
@@ -454,10 +497,14 @@ static inline bool nvshmemSupportsExternalGraphCapture(const cudecompGridDesc_t
 }
 
 static inline bool isManagedPointer(void* ptr) {
+#if !CUDECOMP_BUILD_CPU_ONLY
   // Check if input pointer is managed
   cudaPointerAttributes attr;
   CHECK_CUDA(cudaPointerGetAttributes(&attr, ptr));
   return attr.type == cudaMemoryTypeManaged;
+#else
+  return false;
+#endif
 }
 
 static void setCommInfo(cudecompHandle_t& handle, cudecompGridDesc_t& grid_desc, MPI_Comm mpi_comm,
@@ -647,6 +694,7 @@ template <typename T> std::unordered_map<T, unsigned int> getUniqueIds(const std
   return ids;
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 // Custom deleter for NCCL communicators
 struct ncclCommDeleter {
   void operator()(ncclComm_t* comm) {
@@ -662,6 +710,7 @@ static inline ncclComm createNcclComm(ncclComm_t comm) {
   return std::shared_ptr<ncclComm_t>(new ncclComm_t(comm), ncclCommDeleter());
 }
 
+#endif
 // Helper to check if pencil axis has empty pencils
 static inline bool checkForEmptyPencils(const cudecompGridDesc_t grid_desc, int axis) {
   int j = 0;
diff --git a/include/internal/cudecomp_kernels.h b/include/internal/cudecomp_kernels.h
index dfe041a..132b6cd 100644
--- a/include/internal/cudecomp_kernels.h
+++ b/include/internal/cudecomp_kernels.h
@@ -21,9 +21,13 @@
 #include <cstddef>
 #include <cstdint>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda/std/complex>
+#endif
 
 #include "internal/common.h"
+#include "internal/utils.h"
 
 namespace cudecomp {
 
@@ -58,9 +62,9 @@ template <typename T> struct cudecompNvshmemP2PParams {
 void cudecomp_nvshmem_alltoallv(const cudecompNvshmemA2AParams<float>& params, uint64_t* sig_addr, cudaStream_t stream);
 void cudecomp_nvshmem_alltoallv(const cudecompNvshmemA2AParams<double>& params, uint64_t* sig_addr,
                                 cudaStream_t stream);
-void cudecomp_nvshmem_alltoallv(const cudecompNvshmemA2AParams<cuda::std::complex<float>>& params, uint64_t* sig_addr,
+void cudecomp_nvshmem_alltoallv(const cudecompNvshmemA2AParams<cudecomp::complex<float>>& params, uint64_t* sig_addr,
                                 cudaStream_t stream);
-void cudecomp_nvshmem_alltoallv(const cudecompNvshmemA2AParams<cuda::std::complex<double>>& params, uint64_t* sig_addr,
+void cudecomp_nvshmem_alltoallv(const cudecompNvshmemA2AParams<cudecomp::complex<double>>& params, uint64_t* sig_addr,
                                 cudaStream_t stream);
 
 void cudecomp_nvshmem_alltoallv_p2p(cudecompHandle_t handle, const cudecompNvshmemP2PParams<float>& params,
@@ -68,10 +72,10 @@ void cudecomp_nvshmem_alltoallv_p2p(cudecompHandle_t handle, const cudecompNvshm
 void cudecomp_nvshmem_alltoallv_p2p(cudecompHandle_t handle, const cudecompNvshmemP2PParams<double>& params,
                                     uint64_t* sig_addr, cudaStream_t stream);
 void cudecomp_nvshmem_alltoallv_p2p(cudecompHandle_t handle,
-                                    const cudecompNvshmemP2PParams<cuda::std::complex<float>>& params,
+                                    const cudecompNvshmemP2PParams<cudecomp::complex<float>>& params,
                                     uint64_t* sig_addr, cudaStream_t stream);
 void cudecomp_nvshmem_alltoallv_p2p(cudecompHandle_t handle,
-                                    const cudecompNvshmemP2PParams<cuda::std::complex<double>>& params,
+                                    const cudecompNvshmemP2PParams<cudecomp::complex<double>>& params,
                                     uint64_t* sig_addr, cudaStream_t stream);
 #endif
 
@@ -90,10 +94,10 @@ void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle, cudecompBatchedD2DM
 void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle, cudecompBatchedD2DMemcpy3DParams<double>& params,
                                     cudaStream_t stream);
 void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle,
-                                    cudecompBatchedD2DMemcpy3DParams<cuda::std::complex<float>>& params,
+                                    cudecompBatchedD2DMemcpy3DParams<cudecomp::complex<float>>& params,
                                     cudaStream_t stream);
 void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle,
-                                    cudecompBatchedD2DMemcpy3DParams<cuda::std::complex<double>>& params,
+                                    cudecompBatchedD2DMemcpy3DParams<cudecomp::complex<double>>& params,
                                     cudaStream_t stream);
 
 } // namespace cudecomp
diff --git a/include/internal/halo.h b/include/internal/halo.h
index f604657..5ce6976 100644
--- a/include/internal/halo.h
+++ b/include/internal/halo.h
@@ -25,14 +25,19 @@
 #include <sstream>
 #include <vector>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda_runtime.h>
+#endif
 #include <mpi.h>
 
 #include "internal/checks.h"
 #include "internal/comm_routines.h"
 #include "internal/cudecomp_kernels.h"
 #include "internal/nvtx.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include "internal/performance.h"
+#endif
 #include "internal/utils.h"
 
 namespace cudecomp {
@@ -71,6 +76,7 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
   if (halo_extents[dim] == 0) { return; }
 
   cudecompHaloPerformanceSample* current_sample = nullptr;
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable) {
     auto& samples =
         getOrCreateHaloPerformanceSamples(handle, grid_desc,
@@ -83,6 +89,7 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
     // Record start event
     CHECK_CUDA(cudaEventRecord(current_sample->halo_start_event, stream));
   }
+#endif
 
   int count = 0;
   for (int i = 0; i < 3; ++i) {
@@ -108,6 +115,7 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
     c = 0;
   } else if (neighbors[0] == -1 && neighbors[1] == -1) {
     // Single rank in this dimension and not periodic. Return.
+#if !CUDECOMP_BUILD_CPU_ONLY
     if (handle->performance_report_enable && current_sample) {
       // Record end event and advance sample even for early return
       CHECK_CUDA(cudaEventRecord(current_sample->halo_end_event, stream));
@@ -115,6 +123,7 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
                                    createHaloConfig(ax, dim, input, halo_extents.data(), halo_periods.data(),
                                                     padding.data(), getCudecompDataType<T>()));
     }
+#endif
     return;
   } else {
     // For multi-rank cases, check if halos include ranks other than nearest neighbor process (unsupported currently).
@@ -230,12 +239,14 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
     std::array<size_t, 2> offsets{};
     offsets[1] = static_cast<size_t>(alignCountToBytes(halo_size, CUDECOMP_WORKSPACE_ALIGN_BYTES));
 
+#if !CUDECOMP_BUILD_CPU_ONLY
     if (handle->performance_report_enable && current_sample) {
       current_sample->sendrecv_bytes = 0;
       for (int i = 0; i < 2; ++i) {
         if (neighbors[i] != -1) { current_sample->sendrecv_bytes += halo_size * sizeof(T); }
       }
     }
+#endif
     cudecompSendRecvPair(handle, grid_desc, neighbors, send_buff, counts, offsets, recv_buff, counts, offsets, stream,
                          current_sample);
 
@@ -294,17 +305,20 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
     lx[dim] = shape_g_h_p[dim] - halo_extents[dim];
     recv_offsets[1] = getPencilPtrOffset(pinfo_h, lx);
 
+#if !CUDECOMP_BUILD_CPU_ONLY
     if (handle->performance_report_enable && current_sample) {
       current_sample->sendrecv_bytes = 0;
       for (int i = 0; i < 2; ++i) {
         if (neighbors[i] != -1) { current_sample->sendrecv_bytes += halo_size * sizeof(T); }
       }
     }
+#endif
     cudecompSendRecvPair(handle, grid_desc, neighbors, input, counts, send_offsets, input, counts, recv_offsets, stream,
                          current_sample);
   } break;
   }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable && current_sample) {
     // Record end event
     CHECK_CUDA(cudaEventRecord(current_sample->halo_end_event, stream));
@@ -312,6 +326,7 @@ void cudecompUpdateHalos_(int ax, const cudecompHandle_t handle, const cudecompG
                                  createHaloConfig(ax, dim, input, halo_extents.data(), halo_periods.data(),
                                                   padding.data(), getCudecompDataType<T>()));
   }
+#endif
 }
 
 template <typename T>
diff --git a/include/internal/transpose.h b/include/internal/transpose.h
index d5a39dd..d0b22f1 100644
--- a/include/internal/transpose.h
+++ b/include/internal/transpose.h
@@ -25,17 +25,22 @@
 #include <limits>
 #include <vector>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda/std/complex>
 #include <cuda_runtime.h>
 #include <cutensor.h>
+#endif
 #include <mpi.h>
 
 #include "internal/checks.h"
 #include "internal/comm_routines.h"
 #include "internal/cudecomp_kernels.h"
 #include "internal/nvtx.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include "internal/performance.h"
 #include "internal/raii_wrappers.h"
+#endif
 #include "internal/utils.h"
 
 namespace cudecomp {
@@ -48,6 +53,28 @@ static inline bool isTransposeCommPipelined(cudecompTransposeCommBackend_t commT
           commType == CUDECOMP_TRANSPOSE_COMM_MPI_P2P_PL);
 }
 
+#if CUDECOMP_BUILD_CPU_ONLY
+// Host permutation uses the same element strides and axis order as cuTENSOR.
+template <typename T>
+static void localPermute(const cudecompHandle_t, const std::array<int64_t, 3>& extents,
+                         const std::array<int32_t, 3>& order, std::array<int64_t, 3> strides_in,
+                         std::array<int64_t, 3> strides_out, const T* input, T* output, cudaStream_t) {
+  if (!anyNonzeros(strides_in)) { strides_in = {1, extents[0], extents[0] * extents[1]}; }
+  if (!anyNonzeros(strides_out)) { strides_out = {1, extents[order[0]], extents[order[0]] * extents[order[1]]}; }
+  std::array<int64_t, 3> dst_strides{};
+  for (int i = 0; i < 3; ++i) {
+    dst_strides[order[i]] = strides_out[i];
+  }
+  for (int64_t k = 0; k < extents[2]; ++k) {
+    for (int64_t j = 0; j < extents[1]; ++j) {
+      for (int64_t i = 0; i < extents[0]; ++i) {
+        output[i * dst_strides[0] + j * dst_strides[1] + k * dst_strides[2]] =
+            input[i * strides_in[0] + j * strides_in[1] + k * strides_in[2]];
+      }
+    }
+  }
+}
+#else
 #if CUTENSOR_MAJOR >= 2
 static inline cutensorDataType_t getCutensorDataType(float) { return CUTENSOR_R_32F; }
 static inline cutensorDataType_t getCutensorDataType(double) { return CUTENSOR_R_64F; }
@@ -192,6 +219,7 @@ static void localPermute(const cudecompHandle_t handle, const std::array<int64_t
                                      &desc_out, order_out.data(), cuda_type, stream));
 }
 #endif
+#endif
 
 template <typename T>
 static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, const cudecompGridDesc_t grid_desc,
@@ -218,10 +246,12 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
   bool output_has_halos_padding = anyNonzeros(output_halo_extents) || anyNonzeros(output_padding);
   bool pipelined = isTransposeCommPipelined(grid_desc->config.transpose_comm_backend);
   int memcpy_limit = pipelined ? 1 : CUDECOMP_BATCHED_D2D_3D_PARAM_CAPACITY;
-  cudaStreamCaptureStatus capture_status;
+#if !CUDECOMP_BUILD_CPU_ONLY
+  cudaStreamCaptureStatus capture_status = cudaStreamCaptureStatusNone;
   CHECK_CUDA(cudaStreamIsCapturing(stream, &capture_status));
   bool use_internal_graphs = handle->cuda_graphs_enable && capture_status == cudaStreamCaptureStatusNone;
 
+#endif
   // Set axis values
   int ax_a = ax;
   int ax_b = (fwd ? ax_a + 1 : ax_a + 2) % 3;
@@ -308,6 +338,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
 #endif
 
   cudecompTransposePerformanceSample* current_sample = nullptr;
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable) {
     auto& samples =
         getOrCreateTransposePerformanceSamples(handle, grid_desc,
@@ -322,6 +353,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
     // Record start event
     CHECK_CUDA(cudaEventRecord(current_sample->transpose_start_event, stream));
   }
+#endif
 
   // Adjust pointers to handle special cases
   bool direct_pack = false;
@@ -332,6 +364,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
       if (inplace) {
         if (halos_padding_equal) {
           // Single rank, in place, Pack -> Unpack: No transpose necessary.
+#if !CUDECOMP_BUILD_CPU_ONLY
           if (handle->performance_report_enable) {
             // Record performance data
             CHECK_CUDA(cudaEventRecord(current_sample->transpose_end_event, stream));
@@ -340,6 +373,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
                                                                     output_halo_extents.data(), input_padding.data(),
                                                                     output_padding.data(), getCudecompDataType<T>()));
           }
+#endif
           return;
         }
       } else {
@@ -385,7 +419,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
     } else if (transposeBackendRequiresMpi(grid_desc->config.transpose_comm_backend)) {
       // Note: For MPI, disable special cases if input or output pointers are to managed memory
       // since MPI performance directly from managed memory is not great
-      if (isManagedPointer(input) || isManagedPointer(output)) { enable = false; }
+      if ((isManagedPointer(input) || isManagedPointer(output))) { enable = false; }
 
       // Note: For MPI, disable special cases if communicator has an MNNVL connection and the workspace
       // is fabric allocated. This forces MPI comms to always use the fabric allocated workspace
@@ -462,13 +496,18 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
         auto dtype = getCudecompDataType<T>();
         auto key = std::tie(i1, o1, ax, dir, pinfo_a_h, pinfo_b_h, dtype);
 
+#if !CUDECOMP_BUILD_CPU_ONLY
         if (use_internal_graphs && grid_desc->graph_cache.cached(key)) {
           grid_desc->graph_cache.replay(key, stream);
-        } else {
+        } else
+#endif
+        {
           cudaStream_t graph_stream = stream;
+#if !CUDECOMP_BUILD_CPU_ONLY
           if (use_internal_graphs && splits_a.size() > 1) {
             graph_stream = grid_desc->graph_cache.startCapture(key, stream);
           }
+#endif
 
           for (int j = 1; j < splits_a.size() + 1; ++j) {
             int src_rank, dst_rank;
@@ -503,19 +542,23 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
             }
 
             localPermute(handle, extents, order, strides_in, strides_out, src, dst, graph_stream);
+#if !CUDECOMP_BUILD_CPU_ONLY
 #if CUDART_VERSION >= 11010
             CHECK_CUDA(cudaEventRecordWithFlags(grid_desc->events[dst_rank], graph_stream,
                                                 use_internal_graphs && splits_a.size() > 1 ? cudaEventRecordExternal
                                                                                            : cudaEventRecordDefault));
 #else
-            CHECK_CUDA(cudaEventRecord((grid_desc->events[dst_rank], graph_stream));
+            CHECK_CUDA(cudaEventRecord(grid_desc->events[dst_rank], graph_stream));
+#endif
 #endif
           }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
           if (use_internal_graphs && splits_a.size() > 1) {
             grid_desc->graph_cache.endCapture(key);
             grid_desc->graph_cache.replay(key, stream);
           }
+#endif
         }
       } else {
         T* src = i1 + getPencilPtrOffset(pinfo_a_h, input_halo_extents);
@@ -538,13 +581,18 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
       auto dtype = getCudecompDataType<T>();
       auto key = std::tie(i1, o1, ax, dir, pinfo_a_h, pinfo_b_h, dtype);
 
+#if !CUDECOMP_BUILD_CPU_ONLY
       if (use_internal_graphs && grid_desc->graph_cache.cached(key)) {
         grid_desc->graph_cache.replay(key, stream);
-      } else {
+      } else
+#endif
+      {
         cudaStream_t graph_stream = stream;
+#if !CUDECOMP_BUILD_CPU_ONLY
         if (use_internal_graphs && pipelined && splits_a.size() > 1) {
           graph_stream = grid_desc->graph_cache.startCapture(key, stream);
         }
+#endif
 
         for (int j = 1; j < splits_a.size() + 1; ++j) {
           int src_rank, dst_rank;
@@ -596,6 +644,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
             cudecomp_batched_d2d_memcpy_3d(handle, memcpy_params, graph_stream);
             memcpy_count = 0;
           }
+#if !CUDECOMP_BUILD_CPU_ONLY
 #if CUDART_VERSION >= 11010
           if (pipelined) {
             CHECK_CUDA(cudaEventRecordWithFlags(grid_desc->events[dst_rank], graph_stream,
@@ -603,18 +652,22 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
                                                                                            : cudaEventRecordDefault));
           }
 #else
-          if (pipelined) CHECK_CUDA(cudaEventRecord((grid_desc->events[dst_rank], graph_stream));
+          if (pipelined) CHECK_CUDA(cudaEventRecord(grid_desc->events[dst_rank], graph_stream));
+#endif
 #endif
         }
+#if !CUDECOMP_BUILD_CPU_ONLY
         if (use_internal_graphs && pipelined && splits_a.size() > 1) {
           grid_desc->graph_cache.endCapture(key);
           grid_desc->graph_cache.replay(key, stream);
         }
+#endif
       }
     }
 
     if (o1 == output) {
       // o1 is output. Return.
+#if !CUDECOMP_BUILD_CPU_ONLY
       if (handle->performance_report_enable) {
         CHECK_CUDA(cudaEventRecord(current_sample->transpose_end_event, stream));
         advanceTransposePerformanceSample(handle, grid_desc,
@@ -622,16 +675,19 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
                                                                 output_halo_extents.data(), input_padding.data(),
                                                                 output_padding.data(), getCudecompDataType<T>()));
       }
+#endif
       return;
     }
   } else {
     // For special cases that skip packing and are pipelined, need to record events to
     // enforce input data dependency
+#if !CUDECOMP_BUILD_CPU_ONLY
     if (pipelined) {
       for (int j = 0; j < splits_a.size(); ++j) {
         CHECK_CUDA(cudaEventRecord(grid_desc->events[j], stream));
       }
     }
+#endif
   }
 
   // Communicate
@@ -891,6 +947,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
     }
   }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->performance_report_enable) {
     // Record performance data
     CHECK_CUDA(cudaEventRecord(current_sample->transpose_end_event, stream));
@@ -899,6 +956,7 @@ static void cudecompTranspose_(int ax, int dir, const cudecompHandle_t handle, c
                                                             output_halo_extents.data(), input_padding.data(),
                                                             output_padding.data(), getCudecompDataType<T>()));
   }
+#endif
 }
 
 template <typename T>
diff --git a/include/internal/utils.h b/include/internal/utils.h
index 61578d5..c094da2 100644
--- a/include/internal/utils.h
+++ b/include/internal/utils.h
@@ -18,10 +18,23 @@
 #ifndef CUDECOMP_UTILS_H
 #define CUDECOMP_UTILS_H
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda/std/complex>
+#else
+#include <complex>
+#endif
 
 #include "cudecomp.h"
 
+namespace cudecomp {
+#if !CUDECOMP_BUILD_CPU_ONLY
+template <typename T> using complex = cuda::std::complex<T>;
+#else
+template <typename T> using complex = std::complex<T>;
+#endif
+} // namespace cudecomp
+
 inline bool operator==(const cudecompPencilInfo_t& a, const cudecompPencilInfo_t& b) {
   if (a.size != b.size) return false;
   for (int i = 0; i < 3; ++i) {
@@ -35,8 +48,8 @@ inline bool operator==(const cudecompPencilInfo_t& a, const cudecompPencilInfo_t
 
 inline cudecompDataType_t getCudecompDataType(float) { return CUDECOMP_FLOAT; }
 inline cudecompDataType_t getCudecompDataType(double) { return CUDECOMP_DOUBLE; }
-inline cudecompDataType_t getCudecompDataType(cuda::std::complex<float>) { return CUDECOMP_FLOAT_COMPLEX; }
-inline cudecompDataType_t getCudecompDataType(cuda::std::complex<double>) { return CUDECOMP_DOUBLE_COMPLEX; }
+inline cudecompDataType_t getCudecompDataType(cudecomp::complex<float>) { return CUDECOMP_FLOAT_COMPLEX; }
+inline cudecompDataType_t getCudecompDataType(cudecomp::complex<double>) { return CUDECOMP_DOUBLE_COMPLEX; }
 template <typename T> inline cudecompDataType_t getCudecompDataType() { return getCudecompDataType(T(0)); }
 
 #endif // CUDECOMP_UTILS_H
diff --git a/src/cudecomp.cc b/src/cudecomp.cc
index d7eca15..68db55a 100644
--- a/src/cudecomp.cc
+++ b/src/cudecomp.cc
@@ -33,29 +33,38 @@
 #include <utility>
 #include <vector>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda.h>
 #include <cuda_runtime.h>
+#endif
 #include <mpi.h>
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <nccl.h>
+#endif
 #ifdef ENABLE_NVSHMEM
 #include <nvshmem.h>
 #include <nvshmemx.h>
 #endif
 
 #include "cudecomp.h"
-#include "internal/autotune.h"
 #include "internal/checks.h"
 #include "internal/common.h"
-#include "internal/cuda_wrap.h"
 #include "internal/exceptions.h"
 #include "internal/halo.h"
 #include "internal/hashes.h"
-#include "internal/nvml_wrap.h"
 #include "internal/transpose.h"
+#include "internal/utils.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
+#include "internal/autotune.h"
+#include "internal/cuda_wrap.h"
+#include "internal/nvml_wrap.h"
+#endif
 
 namespace cudecomp {
 namespace {
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 static cudecomp::ncclComm ncclCommFromMPIComm(MPI_Comm mpi_comm) {
   int rank, nranks;
   CHECK_MPI(MPI_Comm_rank(mpi_comm, &rank));
@@ -70,6 +79,7 @@ static cudecomp::ncclComm ncclCommFromMPIComm(MPI_Comm mpi_comm) {
 
   return cudecomp::createNcclComm(nccl_comm);
 }
+#endif
 
 #ifdef ENABLE_NVSHMEM
 static nvshmemProcessState process_nvshmem_state;
@@ -519,6 +529,7 @@ static void gatherGlobalMPIInfo(cudecompHandle_t& handle) {
   CHECK_MPI(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, handle->rank_to_local_rank.data(), 1, MPI_INT,
                           handle->mpi_comm));
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   // Gather MNNVL clique related structures (if supported)
   int dev;
   CHECK_CUDA(cudaGetDevice(&dev));
@@ -592,8 +603,10 @@ static void gatherGlobalMPIInfo(cudecompHandle_t& handle) {
 
   // Copy local rank mapping to clique rank mapping in general case
   if (handle->rank_to_clique_rank.size() == 0) { handle->rank_to_clique_rank = handle->rank_to_local_rank; }
+#endif
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 static void getCudecompEnvVars(cudecompHandle_t& handle) {
   // Check CUDECOMP_ENABLE_NCCL_UBR (NCCL user buffer registration)
   bool nccl_ubr_requested = checkEnvVar("CUDECOMP_ENABLE_NCCL_UBR");
@@ -711,6 +724,7 @@ static void getCudecompEnvVars(cudecompHandle_t& handle) {
   const char* performance_write_dir_str = std::getenv("CUDECOMP_PERFORMANCE_REPORT_WRITE_DIR");
   if (performance_write_dir_str) { handle->performance_report_write_dir = std::string(performance_write_dir_str); }
 }
+#endif
 
 static void resolveRankOrder(cudecompHandle_t handle, cudecompGridDesc_t grid_desc) {
   const char* env_val = std::getenv("CUDECOMP_USE_COL_MAJOR_RANK_ORDER");
@@ -820,11 +834,13 @@ static nvshmemRuntime acquireNvshmemRuntime(cudecompHandle_t& handle) {
 static void releaseUnusedHandleResources(cudecompHandle_t handle, bool release_streams = false) noexcept {
   if (!handle || !handle->initialized) { return; }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   // Grid descriptors retain NCCL communicator references uniformly across each communicator, so collective grid
   // descriptor creation and destruction leave the same reference counts on all of its ranks.
   if (handle->nccl_comm && handle->nccl_comm.use_count() == 1) { handle->nccl_comm.reset(); }
   if (handle->nccl_local_comm && handle->nccl_local_comm.use_count() == 1) { handle->nccl_local_comm.reset(); }
   if (release_streams) { handle->streams.clear(); }
+#endif
 }
 
 static void cleanupFailedGridDescCreate(cudecompHandle_t handle, cudecompGridDesc_t grid_desc,
@@ -885,6 +901,7 @@ static size_t getWorkspaceAllocationSize(cudecompHandle_t handle, size_t request
   return requested_size;
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 static void deregisterNcclBuffer(cudecompHandle_t handle, void* buffer) {
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2, 19, 0)
   auto entry = handle->nccl_ubr_handles.find(buffer);
@@ -898,9 +915,13 @@ static void deregisterNcclBuffer(cudecompHandle_t handle, void* buffer) {
   (void)buffer;
 #endif
 }
+#endif
 
 static void freeOrdinaryWorkspace(cudecompHandle_t handle, void* buffer) {
   if (!buffer) return;
+#if CUDECOMP_BUILD_CPU_ONLY
+  std::free(buffer);
+#else
   deregisterNcclBuffer(handle, buffer);
   if (handle->cuda_cumem_enable) {
 #if CUDART_VERSION >= 11030
@@ -916,6 +937,7 @@ static void freeOrdinaryWorkspace(cudecompHandle_t handle, void* buffer) {
   } else {
     CHECK_CUDA(cudaFree(buffer));
   }
+#endif
 }
 
 #ifdef ENABLE_NVSHMEM
@@ -932,7 +954,9 @@ static void freeNvshmemWorkspace(const nvshmemRuntime& runtime, void* buffer) {
 #endif
 
 static void releaseManagedWorkspaces(cudecompHandle_t handle) {
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (handle->workspace_stream) { CHECK_CUDA(cudaStreamSynchronize(*handle->workspace_stream)); }
+#endif
 
   freeOrdinaryWorkspace(handle, handle->ordinary_workspace.ptr);
   handle->ordinary_workspace.ptr = nullptr;
@@ -947,15 +971,18 @@ static void releaseManagedWorkspaces(cudecompHandle_t handle) {
   handle->nvshmem_workspace.size = 0;
   handle->nvshmem_workspace.capture_frozen = false;
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   handle->workspace_ingress_event.reset();
   handle->workspace_egress_event.reset();
   handle->workspace_stream.reset();
+#endif
 }
 
 // NVSHMEM allocation sizes must be normalized with getWorkspaceAllocationSize before calling this function.
 static void allocateWorkspace(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void** buffer,
                               size_t buffer_size_bytes, bool use_nvshmem);
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 static void registerWorkspaceWithNccl(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void* buffer,
                                       size_t size) {
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2, 19, 0)
@@ -981,7 +1008,6 @@ static void registerWorkspaceWithNccl(cudecompHandle_t handle, cudecompGridDesc_
   (void)size;
 #endif
 }
-
 static bool workspaceRegisteredWithNccl(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void* buffer) {
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2, 19, 0)
   if (!handle->nccl_enable_ubr) return true;
@@ -1000,6 +1026,7 @@ static bool workspaceRegisteredWithNccl(cudecompHandle_t handle, cudecompGridDes
   return true;
 #endif
 }
+#endif
 
 static bool allRanksTrue(cudecompHandle_t handle, bool local_value) {
   int value = local_value;
@@ -1018,7 +1045,9 @@ static void* prepareManagedWorkspace(cudecompHandle_t handle, cudecompGridDesc_t
         "automatic workspace cannot grow after being used by an externally captured CUDA Graph; use a new handle");
   }
   if (workspace.size < required_size) {
+#if !CUDECOMP_BUILD_CPU_ONLY
     CHECK_CUDA(cudaStreamSynchronize(*handle->workspace_stream));
+#endif
     if (workspace.ptr) {
 #ifdef ENABLE_NVSHMEM
       if (use_nvshmem) {
@@ -1039,7 +1068,9 @@ static void* prepareManagedWorkspace(cudecompHandle_t handle, cudecompGridDesc_t
     if (use_nvshmem) { workspace.nvshmem_runtime = grid_desc->nvshmem_runtime; }
 #endif
   } else if (!use_nvshmem) {
+#if !CUDECOMP_BUILD_CPU_ONLY
     registerWorkspaceWithNccl(handle, grid_desc, workspace.ptr, workspace.size);
+#endif
   }
 
   return workspace.ptr;
@@ -1049,6 +1080,14 @@ template <typename Function>
 static void runWithWorkspace(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void* work, size_t required_size,
                              bool use_nvshmem, bool graph_capture_supported, cudaStream_t caller_stream,
                              Function&& function) {
+#if CUDECOMP_BUILD_CPU_ONLY
+  if (work) {
+    function(work, nullptr);
+  } else {
+    std::lock_guard<std::mutex> lock(handle->workspace_mutex);
+    function(prepareManagedWorkspace(handle, grid_desc, required_size, false), nullptr);
+  }
+#else
   cudaStreamCaptureStatus capture_status;
   CHECK_CUDA(cudaStreamIsCapturing(caller_stream, &capture_status));
   if (capture_status == cudaStreamCaptureStatusInvalidated) {
@@ -1106,6 +1145,7 @@ static void runWithWorkspace(cudecompHandle_t handle, cudecompGridDesc_t grid_de
   function(managed_work, managed_stream);
   CHECK_CUDA(cudaEventRecord(*handle->workspace_egress_event, managed_stream));
   CHECK_CUDA(cudaStreamWaitEvent(caller_stream, *handle->workspace_egress_event, 0));
+#endif
 }
 
 } // namespace
@@ -1136,6 +1176,7 @@ cudecompHandle::~cudecompHandle() noexcept {
     cudecomp::releaseManagedWorkspaces(this);
   } catch (...) {}
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2, 19, 0)
   for (auto& entry : nccl_ubr_handles) {
     for (const auto& ubr_handle : entry.second) {
@@ -1144,6 +1185,7 @@ cudecompHandle::~cudecompHandle() noexcept {
   }
 #endif
 
+#endif
   if (mpi_clique_comm != MPI_COMM_NULL) { MPI_Comm_free(&mpi_clique_comm); }
 
   if (mpi_local_comm != MPI_COMM_NULL) { MPI_Comm_free(&mpi_local_comm); }
@@ -1153,7 +1195,9 @@ cudecompHandle::~cudecompHandle() noexcept {
   if (cutensor_plan_pref) { cutensorDestroyPlanPreference(cutensor_plan_pref); }
 #endif
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   if (nvml_initialized) { cudecomp::nvmlFnTable.pfn_nvmlShutdown(); }
+#endif
 }
 
 cudecompResult_t cudecompInit(cudecompHandle_t* handle_in, MPI_Comm mpi_comm) {
@@ -1170,6 +1214,9 @@ cudecompResult_t cudecompInit(cudecompHandle_t* handle_in, MPI_Comm mpi_comm) {
     CHECK_MPI(MPI_Comm_rank(handle->mpi_local_comm, &handle->local_rank));
     CHECK_MPI(MPI_Comm_size(handle->mpi_local_comm, &handle->local_nranks));
 
+#if CUDECOMP_BUILD_CPU_ONLY
+    gatherGlobalMPIInfo(handle);
+#else
     // Load CUDA driver symbols into table
     initCuFunctionTable();
 
@@ -1255,8 +1302,8 @@ cudecompResult_t cudecompInit(cudecompHandle_t* handle_in, MPI_Comm mpi_comm) {
     CHECK_CUDA(cudaDeviceGetAttribute(&handle->device_num_sms, cudaDevAttrMultiProcessorCount, dev));
     CHECK_CUDA(cudaDeviceGetAttribute(&handle->device_max_threads_per_sm, cudaDevAttrMaxThreadsPerMultiProcessor, dev));
 
+#endif
     handle->initialized = true;
-
     *handle_in = handle;
   }
   CUDECOMP_CATCH_C_API_ERRORS(if (handle) { delete handle; })
@@ -1270,9 +1317,11 @@ cudecompResult_t cudecompFinalize(cudecompHandle_t handle) {
 
     releaseManagedWorkspaces(handle);
 
+#if !CUDECOMP_BUILD_CPU_ONLY
     while (!handle->nccl_ubr_handles.empty()) {
       deregisterNcclBuffer(handle, handle->nccl_ubr_handles.begin()->first);
     }
+#endif
 
     delete handle;
   }
@@ -1341,13 +1390,24 @@ cudecompResult_t cudecompGridDescCreateVersioned(cudecompHandle_t handle, cudeco
     checkConfig(handle, config, autotune_transpose_backend, autotune_halo_backend);
 
     bool autotune_pdims = (config->pdims[0] == 0 && config->pdims[1] == 0);
+#if CUDECOMP_BUILD_CPU_ONLY
+    if (autotune_pdims || autotune_transpose_backend || autotune_halo_backend) {
+      THROW_NOT_SUPPORTED("CPU builds require explicit process-grid dimensions and do not support autotuning");
+    }
+    if (!transposeBackendRequiresMpi(config->transpose_comm_backend) ||
+        !haloBackendRequiresMpi(config->halo_comm_backend)) {
+      THROW_NOT_SUPPORTED("CPU builds require MPI backends");
+    }
+#else
     if (autotune_pdims && !options) { THROW_INVALID_USAGE("options argument cannot be null if autotuning pdims"); }
+#endif
 
     grid_desc = new cudecompGridDesc;
     grid_desc->initialized = true;
     grid_desc->handle = handle;
     grid_desc->config = *config;
     resolveRankOrder(handle, grid_desc);
+#if !CUDECOMP_BUILD_CPU_ONLY
     auto comm_backend = grid_desc->config.transpose_comm_backend;
     auto halo_comm_backend = grid_desc->config.halo_comm_backend;
 
@@ -1370,6 +1430,7 @@ cudecompResult_t cudecompGridDescCreateVersioned(cudecompHandle_t handle, cudeco
                     transposeBackendRequiresNvshmem) ||
         std::any_of(autotune_halo_candidates.begin(), autotune_halo_candidates.end(), haloBackendRequiresNvshmem);
 
+#endif
     // If transpose_mem_order not used, set based on transpose_axis_contiguous settings)
     grid_desc->transpose_mem_order_set = (config->transpose_mem_order[0][0] >= 0);
 
@@ -1402,6 +1463,7 @@ cudecompResult_t cudecompGridDescCreateVersioned(cudecompHandle_t handle, cudeco
       grid_desc->config.gdims_dist[2] = grid_desc->config.gdims[2];
     }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
     // Initialize NCCL communicator if needed
     if (need_nccl) {
       if (!handle->nccl_comm) { handle->nccl_comm = ncclCommFromMPIComm(handle->mpi_comm); }
@@ -1463,6 +1525,7 @@ cudecompResult_t cudecompGridDescCreateVersioned(cudecompHandle_t handle, cudeco
       }
     }
 
+#endif
     // Setup final row and column communicators
     bool final_need_nvshmem = transposeBackendRequiresNvshmem(grid_desc->config.transpose_comm_backend) ||
                               haloBackendRequiresNvshmem(grid_desc->config.halo_comm_backend);
@@ -1480,6 +1543,7 @@ cudecompResult_t cudecompGridDescCreateVersioned(cudecompHandle_t handle, cudeco
     }
 
 #endif
+#if !CUDECOMP_BUILD_CPU_ONLY
     if (transposeBackendRequiresNccl(grid_desc->config.transpose_comm_backend) ||
         haloBackendRequiresNccl(grid_desc->config.halo_comm_backend)) {
       int need_local_nccl_comm =
@@ -1495,6 +1559,7 @@ cudecompResult_t cudecompGridDescCreateVersioned(cudecompHandle_t handle, cudeco
       grid_desc->nccl_local_comm.reset();
     }
 
+#endif
     releaseUnusedHandleResources(handle);
 
     *grid_desc_in = grid_desc;
@@ -1528,7 +1593,9 @@ cudecompResult_t cudecompGridDescDestroy(cudecompHandle_t handle, cudecompGridDe
     checkGridDesc(handle, grid_desc);
 
     // Print performance report if enabled
+#if !CUDECOMP_BUILD_CPU_ONLY
     if (handle->performance_report_enable) { printPerformanceReport(handle, grid_desc); }
+#endif
 
     delete grid_desc;
     grid_desc = nullptr;
@@ -1715,6 +1782,10 @@ namespace cudecomp {
 namespace {
 static void allocateWorkspace(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void** buffer,
                               size_t buffer_size_bytes, bool use_nvshmem) {
+#if CUDECOMP_BUILD_CPU_ONLY
+  *buffer = std::malloc(buffer_size_bytes);
+  if (buffer_size_bytes && !*buffer) { THROW_INTERNAL_ERROR("host workspace allocation failed"); }
+#else
   if (use_nvshmem) {
 #ifdef ENABLE_NVSHMEM
     auto nvshmem_runtime = grid_desc->nvshmem_runtime;
@@ -1828,6 +1899,7 @@ static void allocateWorkspace(cudecompHandle_t handle, cudecompGridDesc_t grid_d
       throw;
     }
   }
+#endif
 }
 } // namespace
 } // namespace cudecomp
@@ -1985,15 +2057,15 @@ cudecompResult_t cudecompTransposeXToY(cudecompHandle_t handle, cudecompGridDesc
                               input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompTransposeXToY(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                              reinterpret_cast<cuda::std::complex<float>*>(output),
-                              reinterpret_cast<cuda::std::complex<float>*>(actual_work), input_halo_extents,
+        cudecompTransposeXToY(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                              reinterpret_cast<cudecomp::complex<float>*>(output),
+                              reinterpret_cast<cudecomp::complex<float>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompTransposeXToY(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                              reinterpret_cast<cuda::std::complex<double>*>(output),
-                              reinterpret_cast<cuda::std::complex<double>*>(actual_work), input_halo_extents,
+        cudecompTransposeXToY(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                              reinterpret_cast<cudecomp::complex<double>*>(output),
+                              reinterpret_cast<cudecomp::complex<double>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       }
@@ -2037,15 +2109,15 @@ cudecompResult_t cudecompTransposeYToZ(cudecompHandle_t handle, cudecompGridDesc
                               input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompTransposeYToZ(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                              reinterpret_cast<cuda::std::complex<float>*>(output),
-                              reinterpret_cast<cuda::std::complex<float>*>(actual_work), input_halo_extents,
+        cudecompTransposeYToZ(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                              reinterpret_cast<cudecomp::complex<float>*>(output),
+                              reinterpret_cast<cudecomp::complex<float>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompTransposeYToZ(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                              reinterpret_cast<cuda::std::complex<double>*>(output),
-                              reinterpret_cast<cuda::std::complex<double>*>(actual_work), input_halo_extents,
+        cudecompTransposeYToZ(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                              reinterpret_cast<cudecomp::complex<double>*>(output),
+                              reinterpret_cast<cudecomp::complex<double>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       }
@@ -2089,15 +2161,15 @@ cudecompResult_t cudecompTransposeZToY(cudecompHandle_t handle, cudecompGridDesc
                               input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompTransposeZToY(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                              reinterpret_cast<cuda::std::complex<float>*>(output),
-                              reinterpret_cast<cuda::std::complex<float>*>(actual_work), input_halo_extents,
+        cudecompTransposeZToY(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                              reinterpret_cast<cudecomp::complex<float>*>(output),
+                              reinterpret_cast<cudecomp::complex<float>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompTransposeZToY(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                              reinterpret_cast<cuda::std::complex<double>*>(output),
-                              reinterpret_cast<cuda::std::complex<double>*>(actual_work), input_halo_extents,
+        cudecompTransposeZToY(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                              reinterpret_cast<cudecomp::complex<double>*>(output),
+                              reinterpret_cast<cudecomp::complex<double>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       }
@@ -2141,15 +2213,15 @@ cudecompResult_t cudecompTransposeYToX(cudecompHandle_t handle, cudecompGridDesc
                               input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompTransposeYToX(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                              reinterpret_cast<cuda::std::complex<float>*>(output),
-                              reinterpret_cast<cuda::std::complex<float>*>(actual_work), input_halo_extents,
+        cudecompTransposeYToX(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                              reinterpret_cast<cudecomp::complex<float>*>(output),
+                              reinterpret_cast<cudecomp::complex<float>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompTransposeYToX(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                              reinterpret_cast<cuda::std::complex<double>*>(output),
-                              reinterpret_cast<cuda::std::complex<double>*>(actual_work), input_halo_extents,
+        cudecompTransposeYToX(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                              reinterpret_cast<cudecomp::complex<double>*>(output),
+                              reinterpret_cast<cudecomp::complex<double>*>(actual_work), input_halo_extents,
                               output_halo_extents, input_padding, output_padding, actual_stream);
         break;
       }
@@ -2196,14 +2268,14 @@ cudecompResult_t cudecompUpdateHalosX(cudecompHandle_t handle, cudecompGridDesc_
                              actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompUpdateHalosX(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                             reinterpret_cast<cuda::std::complex<float>*>(actual_work), halo_extents, halo_periods, dim,
+        cudecompUpdateHalosX(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                             reinterpret_cast<cudecomp::complex<float>*>(actual_work), halo_extents, halo_periods, dim,
                              padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompUpdateHalosX(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                             reinterpret_cast<cuda::std::complex<double>*>(actual_work), halo_extents, halo_periods,
-                             dim, padding, actual_stream);
+        cudecompUpdateHalosX(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                             reinterpret_cast<cudecomp::complex<double>*>(actual_work), halo_extents, halo_periods, dim,
+                             padding, actual_stream);
         break;
       }
     };
@@ -2248,14 +2320,14 @@ cudecompResult_t cudecompUpdateHalosY(cudecompHandle_t handle, cudecompGridDesc_
                              actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompUpdateHalosY(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                             reinterpret_cast<cuda::std::complex<float>*>(actual_work), halo_extents, halo_periods, dim,
+        cudecompUpdateHalosY(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                             reinterpret_cast<cudecomp::complex<float>*>(actual_work), halo_extents, halo_periods, dim,
                              padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompUpdateHalosY(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                             reinterpret_cast<cuda::std::complex<double>*>(actual_work), halo_extents, halo_periods,
-                             dim, padding, actual_stream);
+        cudecompUpdateHalosY(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                             reinterpret_cast<cudecomp::complex<double>*>(actual_work), halo_extents, halo_periods, dim,
+                             padding, actual_stream);
         break;
       }
     };
@@ -2300,14 +2372,14 @@ cudecompResult_t cudecompUpdateHalosZ(cudecompHandle_t handle, cudecompGridDesc_
                              actual_stream);
         break;
       case CUDECOMP_FLOAT_COMPLEX:
-        cudecompUpdateHalosZ(handle, grid_desc, reinterpret_cast<cuda::std::complex<float>*>(input),
-                             reinterpret_cast<cuda::std::complex<float>*>(actual_work), halo_extents, halo_periods, dim,
+        cudecompUpdateHalosZ(handle, grid_desc, reinterpret_cast<cudecomp::complex<float>*>(input),
+                             reinterpret_cast<cudecomp::complex<float>*>(actual_work), halo_extents, halo_periods, dim,
                              padding, actual_stream);
         break;
       case CUDECOMP_DOUBLE_COMPLEX:
-        cudecompUpdateHalosZ(handle, grid_desc, reinterpret_cast<cuda::std::complex<double>*>(input),
-                             reinterpret_cast<cuda::std::complex<double>*>(actual_work), halo_extents, halo_periods,
-                             dim, padding, actual_stream);
+        cudecompUpdateHalosZ(handle, grid_desc, reinterpret_cast<cudecomp::complex<double>*>(input),
+                             reinterpret_cast<cudecomp::complex<double>*>(actual_work), halo_extents, halo_periods, dim,
+                             padding, actual_stream);
         break;
       }
     };
diff --git a/src/cudecomp_kernels_cpu.cc b/src/cudecomp_kernels_cpu.cc
new file mode 100644
index 0000000..378bb63
--- /dev/null
+++ b/src/cudecomp_kernels_cpu.cc
@@ -0,0 +1,57 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cstring>
+
+#include "internal/cudecomp_kernels.h"
+
+namespace cudecomp {
+
+template <typename T>
+static void batchedCopy(cudecompHandle_t handle, cudecompBatchedD2DMemcpy3DParams<T>& params, cudaStream_t stream) {
+  for (int n = 0; n < params.ncopies; ++n) {
+    if (params.extents[2][n] == 0) continue;
+    for (size_t k = 0; k < params.extents[0][n]; ++k) {
+      for (size_t j = 0; j < params.extents[1][n]; ++j) {
+        std::memcpy(params.dest[n] + k * params.dest_strides[0][n] + j * params.dest_strides[1][n],
+                    params.src[n] + k * params.src_strides[0][n] + j * params.src_strides[1][n],
+                    params.extents[2][n] * sizeof(T));
+      }
+    }
+  }
+}
+
+void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle, cudecompBatchedD2DMemcpy3DParams<float>& params,
+                                    cudaStream_t stream) {
+  batchedCopy(handle, params, stream);
+}
+void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle, cudecompBatchedD2DMemcpy3DParams<double>& params,
+                                    cudaStream_t stream) {
+  batchedCopy(handle, params, stream);
+}
+void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle,
+                                    cudecompBatchedD2DMemcpy3DParams<cudecomp::complex<float>>& params,
+                                    cudaStream_t stream) {
+  batchedCopy(handle, params, stream);
+}
+void cudecomp_batched_d2d_memcpy_3d(cudecompHandle_t handle,
+                                    cudecompBatchedD2DMemcpy3DParams<cudecomp::complex<double>>& params,
+                                    cudaStream_t stream) {
+  batchedCopy(handle, params, stream);
+}
+
+} // namespace cudecomp
diff --git a/src/cudecomp_m.cuf b/src/cudecomp_m.cuf
index 1ee1a6e..4e0697c 100644
--- a/src/cudecomp_m.cuf
+++ b/src/cudecomp_m.cuf
@@ -13,6 +13,19 @@
 ! See the License for the specific language governing permissions and
 ! limitations under the License.
 
+! Compile the same wrappers for CUDA device arrays or standard Fortran host arrays.
+! The NVIDIA/GNU directives below accept all supported buffer types without TYPE(*).
+#if CUDECOMP_BUILD_CPU_ONLY
+#define c_devptr c_ptr
+#define c_devloc c_loc
+#define cuda_stream_kind c_intptr_t
+#define CUDECOMP_DEVICE
+#define CUDECOMP_ARRAY real(c_float), target
+#else
+#define CUDECOMP_DEVICE device,
+#define CUDECOMP_ARRAY real(c_float), device
+#endif
+
 module cudecomp_internal
   use, intrinsic :: iso_c_binding
   implicit none
@@ -94,16 +107,21 @@ end module cudecomp_internal
 module cudecomp
   use, intrinsic :: iso_c_binding
   use, intrinsic :: iso_fortran_env, only: int64, real32, real64
+#if !CUDECOMP_BUILD_CPU_ONLY
   use cudafor
+#endif
+  ! Keep public aliases here and import fixed-version types locally in versioned wrappers.
+  ! Marking a duplicate module-level import private can also hide its public alias in GNU Fortran.
   use cudecomp_internal, only: cudecompGridDescConfig => cudecompGridDescConfig_v1, &
                                cudecompGridDescAutotuneOptions => cudecompGridDescAutotuneOptions_v1, &
                                cudecompPencilInfo => cudecompPencilInfo_v1
-  use cudecomp_internal, only: cudecompGridDescConfig_v1, cudecompGridDescAutotuneOptions_v1, &
-                               cudecompPencilInfo_v1
+
+#if !CUDECOMP_BUILD_CPU_ONLY
   private :: cudafor
+#endif
   private :: cudecompGetGridDescConfig_v1, cudecompGridDescAutotuneOptionsSetDefaults_v1, &
-             cudecompGridDescConfigSetDefaults_v1, cudecompGridDescCreate_v1, cudecompGetPencilInfo_v1, &
-             cudecompGridDescConfig_v1, cudecompGridDescAutotuneOptions_v1, cudecompPencilInfo_v1
+             cudecompGridDescConfigSetDefaults_v1, cudecompGridDescCreate_v1, cudecompGetPencilInfo_v1
+  private :: cudecomp_copy_c_string
   ! enumerators
 
   ! enum for cuDecomp transpose backend options
@@ -412,7 +430,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input, output
-      real(c_float), device :: input(*), output(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output
+      CUDECOMP_ARRAY :: input(*), output(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: input_halo_extents(3), output_halo_extents(3)
@@ -431,7 +450,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input, output
-      real(c_float), device :: input(*), output(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output
+      CUDECOMP_ARRAY :: input(*), output(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: input_halo_extents(3), output_halo_extents(3)
@@ -450,7 +470,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input, output
-      real(c_float), device :: input(*), output(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output
+      CUDECOMP_ARRAY :: input(*), output(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: input_halo_extents(3), output_halo_extents(3)
@@ -469,7 +490,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input, output
-      real(c_float), device :: input(*), output(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output
+      CUDECOMP_ARRAY :: input(*), output(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: input_halo_extents(3), output_halo_extents(3)
@@ -488,7 +510,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input
-      real(c_float), device :: input(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input
+      CUDECOMP_ARRAY :: input(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: halo_extents(3)
@@ -508,7 +531,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input
-      real(c_float), device :: input(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input
+      CUDECOMP_ARRAY :: input(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: halo_extents(3)
@@ -528,7 +552,8 @@ module cudecomp
       type(cudecompHandle), value :: handle
       type(cudecompGridDesc), value :: grid_desc
       !dir$ ignore_tkr input
-      real(c_float), device :: input(*)
+      !GCC$ ATTRIBUTES NO_ARG_CHECK :: input
+      CUDECOMP_ARRAY :: input(*)
       type(c_devptr), value :: work
       integer(c_int), value :: dtype
       integer(c_int32_t) :: halo_extents(3)
@@ -578,6 +603,7 @@ contains
 
   ! cudecompGridDesc creation/manipulation functions
   function cudecompGridDescConfigSetDefaults_v1(config) result(res)
+    use cudecomp_internal, only: cudecompGridDescConfig_v1
     implicit none
     type(cudecompGridDescConfig_v1), target :: config
     integer(c_int) :: res
@@ -588,6 +614,7 @@ contains
   end function cudecompGridDescConfigSetDefaults_v1
 
   function cudecompGridDescAutotuneOptionsSetDefaults_v1(options) result(res)
+    use cudecomp_internal, only: cudecompGridDescAutotuneOptions_v1
     implicit none
     type(cudecompGridDescAutotuneOptions_v1), target :: options
     integer(c_int) :: res
@@ -602,6 +629,7 @@ contains
   end function cudecompGridDescAutotuneOptionsSetDefaults_v1
 
   function cudecompGridDescCreate_v1(handle, grid_desc, config, options) result(res)
+    use cudecomp_internal, only: cudecompGridDescConfig_v1, cudecompGridDescAutotuneOptions_v1
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
@@ -633,6 +661,7 @@ contains
 
   ! General functions
   function cudecompGetPencilInfo_v1(handle, grid_desc, pencil_info, axis, halo_extents, padding) result(res)
+    use cudecomp_internal, only: cudecompPencilInfo_v1
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
@@ -659,6 +688,7 @@ contains
   end function cudecompGetPencilInfo_v1
 
   function cudecompGetGridDescConfig_v1(handle, grid_desc, config) result(res)
+    use cudecomp_internal, only: cudecompGridDescConfig_v1
     type(cudecompHandle), value :: handle
     type(cudecompGridDesc), value :: grid_desc
     type(cudecompGridDescConfig_v1), target :: config
@@ -689,7 +719,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    real(real32), pointer, device, contiguous :: buffer(:)
+    real(real32), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(int64) :: buffer_size
     integer(c_int) :: res
 
@@ -703,7 +733,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    real(real64), pointer, device, contiguous :: buffer(:)
+    real(real64), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(int64) :: buffer_size
     integer(c_int) :: res
 
@@ -717,7 +747,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    complex(real32), pointer, device, contiguous :: buffer(:)
+    complex(real32), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(int64) :: buffer_size
     integer(c_int) :: res
 
@@ -731,7 +761,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    complex(real64), pointer, device, contiguous :: buffer(:)
+    complex(real64), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(int64) :: buffer_size
     integer(c_int) :: res
 
@@ -745,7 +775,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    real(real32), pointer, device, contiguous :: buffer(:)
+    real(real32), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(c_int) :: res
 
     type(c_devptr) :: buffer_c
@@ -758,7 +788,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    real(real64), pointer, device, contiguous :: buffer(:)
+    real(real64), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(c_int) :: res
 
     type(c_devptr) :: buffer_c
@@ -771,7 +801,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    complex(real32), pointer, device, contiguous :: buffer(:)
+    complex(real32), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(c_int) :: res
 
     type(c_devptr) :: buffer_c
@@ -784,7 +814,7 @@ contains
     implicit none
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
-    complex(real64), pointer, device, contiguous :: buffer(:)
+    complex(real64), pointer, CUDECOMP_DEVICE contiguous :: buffer(:)
     integer(c_int) :: res
 
     type(c_devptr) :: buffer_c
@@ -801,7 +831,7 @@ contains
     type(c_ptr) :: cstr
 
     cstr = cudecompTransposeCommBackendToStringC(comm_backend)
-    call __cudecomp_copy_c_string(cstr, res)
+    call cudecomp_copy_c_string(cstr, res)
   end function cudecompTransposeCommBackendToString
 
   function cudecompHaloCommBackendToString(comm_backend) result(res)
@@ -812,7 +842,7 @@ contains
     integer(c_int) :: csize
 
     cstr = cudecompHaloCommBackendToStringC(comm_backend)
-    call __cudecomp_copy_c_string(cstr, res)
+    call cudecomp_copy_c_string(cstr, res)
   end function cudecompHaloCommBackendToString
 
   function cudecompGetShiftedRank(handle, grid_desc, axis, dim, displacement, periodic, shifted_rank) &
@@ -838,7 +868,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, output, work
-    real(c_float), device :: input(*), output(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output, work
+    CUDECOMP_ARRAY :: input(*), output(*), work(*)
     integer :: dtype
     integer(cuda_stream_kind), optional :: stream
     integer, optional :: input_halo_extents(3)
@@ -877,7 +908,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, output, work
-    real(c_float), device :: input(*), output(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output, work
+    CUDECOMP_ARRAY :: input(*), output(*), work(*)
     integer :: dtype
     integer(cuda_stream_kind), optional :: stream
     integer, optional :: input_halo_extents(3)
@@ -916,7 +948,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, output, work
-    real(c_float), device :: input(*), output(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output, work
+    CUDECOMP_ARRAY :: input(*), output(*), work(*)
     integer :: dtype
     integer(cuda_stream_kind), optional :: stream
     integer, optional :: input_halo_extents(3)
@@ -955,7 +988,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, output, work
-    real(c_float), device :: input(*), output(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, output, work
+    CUDECOMP_ARRAY :: input(*), output(*), work(*)
     integer :: dtype
     integer(cuda_stream_kind), optional :: stream
     integer, optional :: input_halo_extents(3)
@@ -995,7 +1029,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, work
-    real(c_float), device :: input(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, work
+    CUDECOMP_ARRAY :: input(*), work(*)
     integer :: dtype
     integer :: halo_extents(3)
     logical :: halo_periods(3)
@@ -1028,7 +1063,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, work
-    real(c_float), device :: input(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, work
+    CUDECOMP_ARRAY :: input(*), work(*)
     integer :: dtype
     integer :: halo_extents(3)
     logical :: halo_periods(3)
@@ -1061,7 +1097,8 @@ contains
     type(cudecompHandle) :: handle
     type(cudecompGridDesc) :: grid_desc
     !dir$ ignore_tkr input, work
-    real(c_float), device :: input(*), work(*)
+    !GCC$ ATTRIBUTES NO_ARG_CHECK :: input, work
+    CUDECOMP_ARRAY :: input(*), work(*)
     integer :: dtype
     integer :: halo_extents(3)
     logical :: halo_periods(3)
@@ -1088,7 +1125,7 @@ contains
   end function cudecompUpdateHalosZ
 
   ! Helper function to copy string
-  subroutine __cudecomp_copy_c_string(cstr, fstr)
+  subroutine cudecomp_copy_c_string(cstr, fstr)
     implicit none
     type(c_ptr) :: cstr
     character(len=:), allocatable :: fstr
@@ -1105,6 +1142,6 @@ contains
     else
       fstr = ' '
     endif
-  end subroutine __cudecomp_copy_c_string
+  end subroutine cudecomp_copy_c_string
 
 end module cudecomp
diff --git a/tests/ctest/CMakeLists.txt b/tests/ctest/CMakeLists.txt
index eb15a34..d8b944c 100644
--- a/tests/ctest/CMakeLists.txt
+++ b/tests/ctest/CMakeLists.txt
@@ -45,11 +45,14 @@ configure_file(
 )
 
 add_library(cudecomp_test_support STATIC)
+if (NOT CUDECOMP_BUILD_CPU_ONLY)
+  target_sources(cudecomp_test_support PRIVATE gpu_test_utils.cc)
+  target_link_libraries(cudecomp_test_support PUBLIC NVHPC::CUDART ${NCCL_LIBRARY})
+endif()
 target_sources(cudecomp_test_support
   PRIVATE
   backend_test_context.cc
   backend_utils.cc
-  gpu_test_utils.cc
   mpi_test_utils.cc
   test_utils.cc
 )
@@ -58,6 +61,7 @@ target_include_directories(cudecomp_test_support
   ${CMAKE_CURRENT_SOURCE_DIR}
   ${CMAKE_CURRENT_BINARY_DIR}
   ${CMAKE_CURRENT_SOURCE_DIR}/../../include
+  ${CMAKE_BINARY_DIR}/include
   ${NVHPC_CUDA_INCLUDE_DIR}
   ${NCCL_INCLUDE_DIR}
 )
@@ -78,11 +82,53 @@ endif()
 target_link_libraries(cudecomp_test_support
   PUBLIC
   MPI::MPI_CXX
-  NVHPC::CUDART
-  ${NCCL_LIBRARY}
   ${CUDECOMP_TEST_FRAMEWORK_TARGET}
 )
 
+if (CUDECOMP_BUILD_CPU_ONLY)
+  foreach(kind api transpose halo)
+    if (kind STREQUAL "api")
+      set(source cpu_api_tests.cc)
+    else()
+      set(source ${kind}_tests.cc)
+    endif()
+    add_executable(cudecomp_test_${kind} ${source} mpi_test_main.cc)
+    target_link_libraries(cudecomp_test_${kind} PRIVATE cudecomp cudecomp_test_support)
+    add_test(NAME cudecomp_cpu_${kind}
+      COMMAND ${MPIEXEC_EXECUTABLE} ${CUDECOMP_TEST_MPIEXEC_PREFLAGS_LIST}
+              ${MPIEXEC_NUMPROC_FLAG} 4 $<TARGET_FILE:cudecomp_test_${kind}>)
+    set_tests_properties(cudecomp_cpu_${kind} PROPERTIES
+      TIMEOUT ${CUDECOMP_TEST_TIMEOUT} LABELS "cpu;mpi")
+  endforeach()
+
+  if (CUDECOMP_BUILD_FORTRAN)
+    foreach(kind api transpose halo)
+      if (kind STREQUAL "api")
+        set(source fortran_cpu_api_test.f90)
+      else()
+        set(source fortran_${kind}_tests.f90)
+      endif()
+      add_executable(cudecomp_test_fortran_cpu_${kind} ${source})
+      target_compile_definitions(cudecomp_test_fortran_cpu_${kind} PRIVATE CUDECOMP_BUILD_CPU_ONLY)
+      if (CMAKE_VERSION VERSION_LESS 3.18)
+        target_compile_options(cudecomp_test_fortran_cpu_${kind} PRIVATE -cpp)
+      else()
+        set_target_properties(cudecomp_test_fortran_cpu_${kind} PROPERTIES Fortran_PREPROCESS ON)
+      endif()
+      target_compile_options(cudecomp_test_fortran_cpu_${kind} PRIVATE
+        $<$<Fortran_COMPILER_ID:GNU>:-ffree-line-length-none>)
+      target_include_directories(cudecomp_test_fortran_cpu_${kind} PRIVATE ${CMAKE_BINARY_DIR}/include)
+      target_link_libraries(cudecomp_test_fortran_cpu_${kind} PRIVATE cudecomp_fort cudecomp)
+      add_test(NAME cudecomp_fortran_cpu_${kind}
+        COMMAND ${MPIEXEC_EXECUTABLE} ${CUDECOMP_TEST_MPIEXEC_PREFLAGS_LIST}
+                ${MPIEXEC_NUMPROC_FLAG} 4 $<TARGET_FILE:cudecomp_test_fortran_cpu_${kind}>)
+      set_tests_properties(cudecomp_fortran_cpu_${kind} PROPERTIES
+        TIMEOUT ${CUDECOMP_TEST_TIMEOUT} LABELS "cpu;fortran;mpi")
+    endforeach()
+  endif()
+  return()
+endif()
+
 add_executable(cudecomp_test_api)
 target_sources(cudecomp_test_api
   PRIVATE
diff --git a/tests/ctest/backend_test_context.cc b/tests/ctest/backend_test_context.cc
index 093f780..1b87b67 100644
--- a/tests/ctest/backend_test_context.cc
+++ b/tests/ctest/backend_test_context.cc
@@ -15,6 +15,7 @@
 namespace cudecomp_test {
 namespace {
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 constexpr const char* kKeepaliveBackendEnv = "CUDECOMP_TEST_KEEPALIVE_BACKEND";
 
 class SharedBackendTestContextState {
@@ -60,6 +61,7 @@ std::string requestedKeepaliveBackend() {
   if (backend == "nccl" || backend == "nvshmem") return backend;
   return {};
 }
+#endif
 
 } // namespace
 
@@ -73,6 +75,7 @@ testing::AssertionResult BackendTestContext::initialize(const MpiTestComm& world
   active_comm_ = nullptr;
   handle_ = nullptr;
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   const std::string keepalive_backend = requestedKeepaliveBackend();
   const bool use_shared_context =
       !keepalive_backend.empty() && keepalive_backend == backend_label && active_ranks == world_comm.size();
@@ -110,6 +113,7 @@ testing::AssertionResult BackendTestContext::initialize(const MpiTestComm& world
     handle_ = shared_context->handle();
     return testing::AssertionSuccess();
   }
+#endif
 
   resetSharedBackendTestContext();
 
@@ -121,8 +125,10 @@ testing::AssertionResult BackendTestContext::initialize(const MpiTestComm& world
     return testing::AssertionSuccess();
   }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   *setup_decision = initializeGpuForTest(local_active_comm_, check_nccl);
   if (setup_decision->skip || setup_decision->fail) return testing::AssertionSuccess();
+#endif
 
   const cudecompResult_t init_result = cudecompInit(&local_handle_, local_active_comm_.mpiComm());
   local_handle_guard_ = std::make_unique<cudecompHandleGuard>(local_handle_);
@@ -134,6 +140,10 @@ testing::AssertionResult BackendTestContext::initialize(const MpiTestComm& world
   return testing::AssertionSuccess();
 }
 
-void resetSharedBackendTestContext() { shared_context.reset(); }
+void resetSharedBackendTestContext() {
+#if !CUDECOMP_BUILD_CPU_ONLY
+  shared_context.reset();
+#endif
+}
 
 } // namespace cudecomp_test
diff --git a/tests/ctest/cpu_api_tests.cc b/tests/ctest/cpu_api_tests.cc
new file mode 100644
index 0000000..5ce905a
--- /dev/null
+++ b/tests/ctest/cpu_api_tests.cc
@@ -0,0 +1,69 @@
+/* SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#include <algorithm>
+
+#include "internal/common.h"
+#include "test_utils.h"
+
+TEST(CpuInitTest, HostAllocationAndUnsupportedFeatures) {
+  const auto comm = cudecomp_test::MpiTestComm::world();
+  cudecompHandle_t handle = nullptr;
+  CHECK_CUDECOMP_GLOBAL(comm, cudecompInit(&handle, comm.mpiComm()));
+  cudecomp_test::cudecompHandleGuard handle_guard(handle);
+  EXPECT_FALSE(handle->nvml_initialized);
+  EXPECT_FALSE(handle->performance_report_enable);
+
+  cudecompGridDescConfig_t config;
+  ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescConfigSetDefaults(&config));
+  config.gdims[0] = config.gdims[1] = config.gdims[2] = 8;
+  config.pdims[0] = comm.size();
+  config.pdims[1] = 1;
+  cudecompGridDesc_t grid = nullptr;
+  config.transpose_comm_backend = CUDECOMP_TRANSPOSE_COMM_NCCL;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, nullptr));
+  config.transpose_comm_backend = CUDECOMP_TRANSPOSE_COMM_MPI_P2P;
+  config.halo_comm_backend = CUDECOMP_HALO_COMM_NCCL;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, nullptr));
+  config.halo_comm_backend = CUDECOMP_HALO_COMM_MPI;
+  config.transpose_comm_backend = CUDECOMP_TRANSPOSE_COMM_NVSHMEM;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, nullptr));
+  config.transpose_comm_backend = CUDECOMP_TRANSPOSE_COMM_MPI_P2P;
+  config.halo_comm_backend = CUDECOMP_HALO_COMM_NVSHMEM;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, nullptr));
+  config.halo_comm_backend = CUDECOMP_HALO_COMM_MPI;
+
+  cudecompGridDescAutotuneOptions_t tune;
+  ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&tune));
+  tune.autotune_transpose_backend = true;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, &tune));
+  tune.autotune_transpose_backend = false;
+  tune.autotune_halo_backend = true;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, &tune));
+  tune.autotune_halo_backend = false;
+  config.pdims[0] = config.pdims[1] = 0;
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, nullptr));
+  EXPECT_EQ(CUDECOMP_RESULT_NOT_SUPPORTED, cudecompGridDescCreate(handle, &grid, &config, &tune));
+  config.pdims[0] = comm.size();
+  config.pdims[1] = 1;
+
+  CHECK_CUDECOMP_GLOBAL(comm, cudecompGridDescCreate(handle, &grid, &config, &tune));
+  cudecomp_test::gridDescGuard grid_guard(handle, grid);
+  void* buffer = nullptr;
+  CHECK_CUDECOMP_GLOBAL(comm, cudecompMalloc(handle, grid, &buffer, 512 * sizeof(float)));
+  cudecomp_test::cudecompBufferGuard buffer_guard(handle, grid, buffer);
+  std::fill_n(static_cast<float*>(buffer), 512, 7.0f);
+  EXPECT_EQ(static_cast<float*>(buffer)[511], 7.0f);
+  const int32_t extents[3] = {0, 0, 0};
+  const bool periods[3] = {false, false, false};
+  for (auto halo : {cudecompUpdateHalosX, cudecompUpdateHalosY, cudecompUpdateHalosZ}) {
+    EXPECT_EQ(CUDECOMP_RESULT_SUCCESS, halo(handle, grid, buffer, nullptr, CUDECOMP_FLOAT, extents, periods, 0, nullptr,
+                                            reinterpret_cast<cudaStream_t>(1)));
+  }
+  for (auto transpose : {cudecompTransposeXToY, cudecompTransposeYToZ, cudecompTransposeZToY, cudecompTransposeYToX}) {
+    EXPECT_EQ(CUDECOMP_RESULT_SUCCESS, transpose(handle, grid, buffer, buffer, nullptr, CUDECOMP_FLOAT, nullptr,
+                                                 nullptr, nullptr, nullptr, reinterpret_cast<cudaStream_t>(1)));
+  }
+  EXPECT_TRUE(std::all_of(static_cast<float*>(buffer), static_cast<float*>(buffer) + 512,
+                          [](float value) { return value == 7.0f; }));
+}
diff --git a/tests/ctest/fortran_cpu_api_test.f90 b/tests/ctest/fortran_cpu_api_test.f90
new file mode 100644
index 0000000..db70598
--- /dev/null
+++ b/tests/ctest/fortran_cpu_api_test.f90
@@ -0,0 +1,22 @@
+! SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+! SPDX-License-Identifier: Apache-2.0
+
+program fortran_cpu_api_test
+  use cudecomp
+  use mpi_f08
+  implicit none
+  type(cudecompHandle) :: handle
+  integer :: ierr
+
+  call MPI_Init(ierr)
+  call check(cudecompInit(handle, MPI_COMM_WORLD))
+  call check(cudecompFinalize(handle))
+  call check(cudecompInit(handle, MPI_COMM_WORLD%MPI_VAL))
+  call check(cudecompFinalize(handle))
+  call MPI_Finalize(ierr)
+contains
+  subroutine check(result)
+    integer, intent(in) :: result
+    if (result /= CUDECOMP_RESULT_SUCCESS) call MPI_Abort(MPI_COMM_WORLD, result, ierr)
+  end subroutine
+end program
diff --git a/tests/ctest/fortran_halo_case.inc b/tests/ctest/fortran_halo_case.inc
index 6cab334..d8eaa26 100644
--- a/tests/ctest/fortran_halo_case.inc
+++ b/tests/ctest/fortran_halo_case.inc
@@ -58,8 +58,8 @@ subroutine RUN_HALO_AXIS(scenario, axis, halo_extents, padding, halo_periods, me
   ARRTYPE, allocatable :: initial(:)
   ARRTYPE, allocatable :: expected(:)
   ARRTYPE, allocatable :: actual(:)
-  ARRTYPE, allocatable, device, target :: data_d(:)
-  ARRTYPE, pointer, device, contiguous :: work_d(:)
+  ARRTYPE, allocatable, CUDECOMP_TEST_DEVICE target :: data_d(:)
+  ARRTYPE, pointer, CUDECOMP_TEST_DEVICE contiguous :: work_d(:)
 
   grid_desc_created = .false.
   work_allocated = .false.
diff --git a/tests/ctest/fortran_halo_tests.f90 b/tests/ctest/fortran_halo_tests.f90
index 92adcf4..2ac58d4 100644
--- a/tests/ctest/fortran_halo_tests.f90
+++ b/tests/ctest/fortran_halo_tests.f90
@@ -3,7 +3,12 @@
 
 module cudecomp_fortran_halo_tests
   use, intrinsic :: iso_fortran_env, only: int64, real32, real64
+#if CUDECOMP_BUILD_CPU_ONLY
+#define CUDECOMP_TEST_DEVICE
+#else
   use cudafor
+#define CUDECOMP_TEST_DEVICE device,
+#endif
   use cudecomp
   use mpi
 
@@ -46,7 +51,9 @@ subroutine run_all_tests()
     if (nranks /= halo_test_ranks) then
       call record_failure("Fortran halo test requires exactly 4 MPI ranks")
     else
+#if !CUDECOMP_BUILD_CPU_ONLY
       call initialize_gpu()
+#endif
     endif
 
     if (failures == 0) then
@@ -74,6 +81,7 @@ subroutine run_all_tests()
     if (global_failures /= 0) call exit(1)
   end subroutine run_all_tests
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   subroutine initialize_gpu()
     implicit none
 
@@ -96,6 +104,7 @@ subroutine initialize_gpu()
     status = cudaSetDevice(mod(local_rank, num_devices))
     if (status /= cudaSuccess) call record_failure("cudaSetDevice failed")
   end subroutine initialize_gpu
+#endif
 
   subroutine setup_halo_config(config, axis_contiguous, mem_order, rank_order)
     implicit none
diff --git a/tests/ctest/fortran_transpose_case.inc b/tests/ctest/fortran_transpose_case.inc
index 1173032..ae0fa9d 100644
--- a/tests/ctest/fortran_transpose_case.inc
+++ b/tests/ctest/fortran_transpose_case.inc
@@ -55,11 +55,11 @@ subroutine RUN_TRANSPOSE_SCENARIO(scenario, halo_extents_x, halo_extents_y, halo
   ARRTYPE, allocatable :: yref(:)
   ARRTYPE, allocatable :: zref(:)
   ARRTYPE, allocatable :: actual(:)
-  ARRTYPE, allocatable, device, target :: data_a_d(:)
-  ARRTYPE, allocatable, device, target :: data_b_d(:)
-  ARRTYPE, pointer, device, contiguous :: input_d(:)
-  ARRTYPE, pointer, device, contiguous :: output_d(:)
-  ARRTYPE, pointer, device, contiguous :: work_d(:)
+  ARRTYPE, allocatable, CUDECOMP_TEST_DEVICE target :: data_a_d(:)
+  ARRTYPE, allocatable, CUDECOMP_TEST_DEVICE target :: data_b_d(:)
+  ARRTYPE, pointer, CUDECOMP_TEST_DEVICE contiguous :: input_d(:)
+  ARRTYPE, pointer, CUDECOMP_TEST_DEVICE contiguous :: output_d(:)
+  ARRTYPE, pointer, CUDECOMP_TEST_DEVICE contiguous :: work_d(:)
 
   grid_desc_created = .false.
   work_allocated = .false.
diff --git a/tests/ctest/fortran_transpose_tests.f90 b/tests/ctest/fortran_transpose_tests.f90
index 21fb4bb..1aab4ed 100644
--- a/tests/ctest/fortran_transpose_tests.f90
+++ b/tests/ctest/fortran_transpose_tests.f90
@@ -3,7 +3,12 @@
 
 module cudecomp_fortran_transpose_tests
   use, intrinsic :: iso_fortran_env, only: int64, real32, real64
+#if CUDECOMP_BUILD_CPU_ONLY
+#define CUDECOMP_TEST_DEVICE
+#else
   use cudafor
+#define CUDECOMP_TEST_DEVICE device,
+#endif
   use cudecomp
   use mpi
 
@@ -48,7 +53,9 @@ subroutine run_all_tests()
     if (nranks /= transpose_test_ranks) then
       call record_failure("Fortran transpose test requires exactly 4 MPI ranks")
     else
+#if !CUDECOMP_BUILD_CPU_ONLY
       call initialize_gpu()
+#endif
     endif
 
     if (failures == 0) then
@@ -76,6 +83,7 @@ subroutine run_all_tests()
     if (global_failures /= 0) call exit(1)
   end subroutine run_all_tests
 
+#if !CUDECOMP_BUILD_CPU_ONLY
   subroutine initialize_gpu()
     implicit none
 
@@ -98,6 +106,7 @@ subroutine initialize_gpu()
     status = cudaSetDevice(mod(local_rank, num_devices))
     if (status /= cudaSuccess) call record_failure("cudaSetDevice failed")
   end subroutine initialize_gpu
+#endif
 
   subroutine setup_transpose_config(config, axis_contiguous, mem_order, rank_order)
     implicit none
diff --git a/tests/ctest/halo_tests.cc b/tests/ctest/halo_tests.cc
index 1799a6b..52a01e8 100644
--- a/tests/ctest/halo_tests.cc
+++ b/tests/ctest/halo_tests.cc
@@ -13,7 +13,10 @@
 
 #include <mpi.h>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda_runtime.h>
+#endif
 #include <gtest/gtest.h>
 
 #include "cudecomp.h"
@@ -166,6 +169,7 @@ std::vector<HaloCase> haloCasesForLabel(const char* label) {
   return cases;
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 std::vector<HaloCase> externalCudaGraphHaloCases(const char* label) {
   std::vector<HaloCase> cases;
   for (const auto& backend : cudecomp_test::haloBackends()) {
@@ -177,6 +181,7 @@ std::vector<HaloCase> externalCudaGraphHaloCases(const char* label) {
   }
   return cases;
 }
+#endif
 
 bool isInternal(const cudecompPencilInfo_t& pinfo, const std::array<int64_t, 3>& local) {
   return local[0] >= pinfo.halo_extents[pinfo.order[0]] &&
@@ -375,6 +380,21 @@ template <typename T> void runHaloCase(const HaloCase& test_case, bool capture_e
   const auto initial = initializePencil<T>(pinfo, test_case.gdims);
   const auto expected = initializeReference<T>(pinfo, test_case.gdims, test_case.halo_periods);
 
+#if CUDECOMP_BUILD_CPU_ONLY
+  void* work = nullptr;
+  if (!test_case.automatic_workspace) {
+    CHECK_CUDECOMP_GLOBAL(active_comm, cudecompMalloc(handle, grid_desc, &work, workspace_num_elements * dtype_size));
+  }
+  cudecomp_test::cudecompBufferGuard workspace(handle, grid_desc, work);
+  for (int repeat = 0; repeat < 2; ++repeat) {
+    auto actual = initial;
+    for (int dim = 0; dim < 3; ++dim) {
+      CHECK_CUDECOMP_GLOBAL(active_comm, runHalo(handle, grid_desc, test_case.axis, actual.data(), work,
+                                                 test_case.dtype, pinfo, test_case.halo_periods, dim));
+    }
+    EXPECT_TRUE(pencilMatches(expected, actual, pinfo));
+  }
+#else
   T* data_d = nullptr;
   const cudaError_t data_alloc_result = cudaMalloc(&data_d, pinfo.size * sizeof(*data_d));
   cudecomp_test::cudaBufferGuard data_buffer(data_d);
@@ -460,6 +480,7 @@ template <typename T> void runHaloCase(const HaloCase& test_case, bool capture_e
     if (graph) { CHECK_CUDA_GLOBAL(active_comm, cudaGraphDestroy(graph)); }
     CHECK_CUDA_GLOBAL(active_comm, cudaStreamDestroy(capture_stream));
   }
+#endif
 }
 
 TEST_P(HaloCorrectnessTest, UpdateHalos) {
@@ -473,6 +494,7 @@ TEST_P(HaloCorrectnessTest, UpdateHalos) {
   }
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 TEST_P(ExternalCudaGraphHaloCorrectnessTest, CapturesAndReplaysOperation) {
   const auto test_case = GetParam();
   switch (test_case.dtype) {
@@ -484,7 +506,9 @@ TEST_P(ExternalCudaGraphHaloCorrectnessTest, CapturesAndReplaysOperation) {
   }
 }
 
+#endif
 INSTANTIATE_TEST_SUITE_P(MpiBackends, HaloCorrectnessTest, ::testing::ValuesIn(haloCasesForLabel("mpi")), paramName);
+#if !CUDECOMP_BUILD_CPU_ONLY
 INSTANTIATE_TEST_SUITE_P(NcclBackends, HaloCorrectnessTest, ::testing::ValuesIn(haloCasesForLabel("nccl")), paramName);
 INSTANTIATE_TEST_SUITE_P(NvshmemBackends, HaloCorrectnessTest, ::testing::ValuesIn(haloCasesForLabel("nvshmem")),
                          paramName);
@@ -492,3 +516,5 @@ INSTANTIATE_TEST_SUITE_P(ExternalCudaGraphNcclBackends, ExternalCudaGraphHaloCor
                          ::testing::ValuesIn(externalCudaGraphHaloCases("nccl")), paramName);
 INSTANTIATE_TEST_SUITE_P(ExternalCudaGraphNvshmemBackends, ExternalCudaGraphHaloCorrectnessTest,
                          ::testing::ValuesIn(externalCudaGraphHaloCases("nvshmem")), paramName);
+
+#endif
diff --git a/tests/ctest/mpi_test_main.cc b/tests/ctest/mpi_test_main.cc
index 82f2822..a842c57 100644
--- a/tests/ctest/mpi_test_main.cc
+++ b/tests/ctest/mpi_test_main.cc
@@ -10,6 +10,7 @@
 #include <gtest/gtest.h>
 
 #include "backend_test_context.h"
+#include "cudecomp_config.h"
 #include "gpu_test_utils.h"
 #include "mpi_test_utils.h"
 
@@ -31,12 +32,16 @@ class RankFailurePrinter : public ::testing::EmptyTestEventListener {
 };
 
 bool gpuRequirementsSatisfied(const cudecomp_test::MpiTestComm& world_comm) {
+#if !CUDECOMP_BUILD_CPU_ONLY
   const auto setup_decision = cudecomp_test::checkGpuTestRequirements(world_comm);
   if (!setup_decision.fail) return true;
 
   std::cerr << "[rank " << world_comm.rank() << "] GPU test requirements are not satisfied: " << setup_decision.reason
             << "\n";
   return false;
+#else
+  return true;
+#endif
 }
 
 } // namespace
@@ -48,7 +53,9 @@ int main(int argc, char** argv) {
 
   {
     const auto world_comm = cudecomp_test::MpiTestComm::world();
+#if !CUDECOMP_BUILD_CPU_ONLY
     cudecomp_test::initializeGpuTestRuntime();
+#endif
 
     ::testing::InitGoogleTest(&argc, argv);
     ::testing::UnitTest::GetInstance()->listeners().Append(new RankFailurePrinter(world_comm.rank()));
diff --git a/tests/ctest/test_utils.cc b/tests/ctest/test_utils.cc
index f57947f..5943135 100644
--- a/tests/ctest/test_utils.cc
+++ b/tests/ctest/test_utils.cc
@@ -44,6 +44,7 @@ gridDescGuard::~gridDescGuard() {
   if (grid_desc_) { (void)cudecompGridDescDestroy(handle_, grid_desc_); }
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 cudaBufferGuard::cudaBufferGuard(void* ptr) : ptr_(ptr) {}
 
 cudaBufferGuard::~cudaBufferGuard() { reset(); }
@@ -53,6 +54,7 @@ void cudaBufferGuard::reset(void* ptr) {
   ptr_ = ptr;
 }
 
+#endif
 cudecompBufferGuard::cudecompBufferGuard(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void* ptr)
     : handle_(handle), grid_desc_(grid_desc), ptr_(ptr) {}
 
@@ -62,9 +64,11 @@ cudecompBufferGuard::~cudecompBufferGuard() {
 
 void cudecompBufferGuard::release() noexcept { ptr_ = nullptr; }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 testing::AssertionResult checkCudaGlobal(const MpiTestComm& comm, cudaError_t result, const char* file, int line) {
   return checkLocalFailureGlobal(comm, result == cudaSuccess, "CUDA", cudaGetErrorString(result), file, line);
 }
+#endif
 
 testing::AssertionResult checkCudecompGlobal(const MpiTestComm& comm, cudecompResult_t result, const char* file,
                                              int line) {
diff --git a/tests/ctest/test_utils.h b/tests/ctest/test_utils.h
index e910f50..72f8ce9 100644
--- a/tests/ctest/test_utils.h
+++ b/tests/ctest/test_utils.h
@@ -8,7 +8,10 @@
 
 #include <mpi.h>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda_runtime.h>
+#endif
 #include <gtest/gtest.h>
 
 #include "cudecomp.h"
@@ -39,6 +42,7 @@ class gridDescGuard {
   cudecompGridDesc_t grid_desc_ = nullptr;
 };
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 class cudaBufferGuard {
 public:
   explicit cudaBufferGuard(void* ptr = nullptr);
@@ -52,6 +56,7 @@ class cudaBufferGuard {
   void* ptr_ = nullptr;
 };
 
+#endif
 class cudecompBufferGuard {
 public:
   cudecompBufferGuard(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, void* ptr);
@@ -67,7 +72,9 @@ class cudecompBufferGuard {
   void* ptr_ = nullptr;
 };
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 testing::AssertionResult checkCudaGlobal(const MpiTestComm& comm, cudaError_t result, const char* file, int line);
+#endif
 testing::AssertionResult checkCudecompGlobal(const MpiTestComm& comm, cudecompResult_t result, const char* file,
                                              int line);
 testing::AssertionResult checkMpiGlobal(const MpiTestComm& comm, int result, const char* file, int line);
diff --git a/tests/ctest/transpose_tests.cc b/tests/ctest/transpose_tests.cc
index a2ae3ec..105293e 100644
--- a/tests/ctest/transpose_tests.cc
+++ b/tests/ctest/transpose_tests.cc
@@ -14,7 +14,10 @@
 
 #include <mpi.h>
 
+#include "cudecomp_config.h"
+#if !CUDECOMP_BUILD_CPU_ONLY
 #include <cuda_runtime.h>
+#endif
 #include <gtest/gtest.h>
 
 #include "cudecomp.h"
@@ -305,6 +308,7 @@ std::vector<TransposeCase> transposeCasesForLabel(const char* label) {
   return cases;
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 std::vector<TransposeCase> cudaGraphTransposeCases() {
   std::vector<TransposeCase> cases;
   for (const auto& backend : cudecomp_test::transposeBackends()) {
@@ -344,6 +348,7 @@ std::vector<TransposeCase> ncclUserBufferRegistrationCases() {
   }
   return cases;
 }
+#endif
 
 bool isInternal(const cudecompPencilInfo_t& pinfo, const std::array<int64_t, 3>& local) {
   return local[0] >= pinfo.halo_extents[pinfo.order[0]] &&
@@ -433,6 +438,7 @@ cudecompResult_t runTranspose(cudecompHandle_t handle, cudecompGridDesc_t grid_d
   return CUDECOMP_RESULT_INVALID_USAGE;
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 template <typename T>
 void runAndVerifyTranspose(const cudecomp_test::MpiTestComm& active_comm, cudecompHandle_t handle,
                            cudecompGridDesc_t grid_desc, const TransposeCase& test_case, T* input_d, T* output_d,
@@ -456,6 +462,7 @@ void runAndVerifyTranspose(const cudecomp_test::MpiTestComm& active_comm, cudeco
   EXPECT_TRUE(pencilMatches(output_ref, output, output_info));
 }
 
+#endif
 cudecomp::cudecompCommAxis communicationAxis(TransposeOperation operation) {
   const int ax_a = inputAxis(operation);
   const int ax_b = outputAxis(operation);
@@ -540,6 +547,7 @@ class CudaGraphTransposeCorrectnessTest : public ::testing::TestWithParam<Transp
 class ExternalCudaGraphTransposeCorrectnessTest : public ::testing::TestWithParam<TransposeCase> {};
 class NcclUserBufferRegistrationTest : public ::testing::TestWithParam<TransposeCase> {};
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 TEST(NcclCommunicatorLifecycleTest, RetainsLocalCommunicatorUntilAllDescriptorsReleaseIt) {
   const auto world_comm = cudecomp_test::MpiTestComm::world();
   if (world_comm.size() != 4) { GTEST_SKIP() << "NCCL communicator lifecycle test requires exactly four ranks"; }
@@ -599,6 +607,7 @@ testing::AssertionResult ncclUserBufferRegistrationIsActive(cudecompHandle_t han
 #endif
 }
 
+#endif
 template <typename T>
 void runTransposeCase(const TransposeCase& test_case, GraphExecution graph_execution = GraphExecution::Eager,
                       bool check_nccl_user_buffer_registration = false) {
@@ -673,6 +682,22 @@ void runTransposeCase(const TransposeCase& test_case, GraphExecution graph_execu
   const auto output_ref = initializePencil<T>(output_info, test_case.gdims);
   const int64_t data_num_elements = std::max(input_info.size, output_info.size);
 
+#if CUDECOMP_BUILD_CPU_ONLY
+  std::vector<T> input(data_num_elements), output(data_num_elements);
+  void* work = nullptr;
+  if (!test_case.automatic_workspace) {
+    CHECK_CUDECOMP_GLOBAL(active_comm, cudecompMalloc(handle, grid_desc, &work, workspace_num_elements * dtype_size));
+  }
+  cudecomp_test::cudecompBufferGuard workspace(handle, grid_desc, work);
+  for (int repeat = 0; repeat < 2; ++repeat) {
+    std::copy(input_ref.begin(), input_ref.end(), input.begin());
+    T* dst = test_case.out_of_place ? output.data() : input.data();
+    CHECK_CUDECOMP_GLOBAL(active_comm, runTranspose(handle, grid_desc, test_case.operation, input.data(), dst, work,
+                                                    test_case.dtype, input_info, output_info));
+    std::vector<T> actual(dst, dst + output_ref.size());
+    EXPECT_TRUE(pencilMatches(output_ref, actual, output_info));
+  }
+#else
   T* input_d = nullptr;
   const cudaError_t input_alloc_result = cudaMalloc(&input_d, data_num_elements * sizeof(*input_d));
   cudecomp_test::cudaBufferGuard input_buffer(input_d);
@@ -799,6 +824,7 @@ void runTransposeCase(const TransposeCase& test_case, GraphExecution graph_execu
     EXPECT_EQ(handle->nccl_ubr_handles.count(work_d), 0);
 #endif
   }
+#endif
 }
 
 TEST_P(TransposeCorrectnessTest, DirectOperation) {
@@ -812,6 +838,7 @@ TEST_P(TransposeCorrectnessTest, DirectOperation) {
   }
 }
 
+#if !CUDECOMP_BUILD_CPU_ONLY
 TEST_P(CudaGraphTransposeCorrectnessTest, CapturesAndReplaysPackingGraph) {
   const auto test_case = GetParam();
   switch (test_case.dtype) {
@@ -849,8 +876,10 @@ TEST_P(NcclUserBufferRegistrationTest, DirectOperation) {
 #endif
 }
 
+#endif
 INSTANTIATE_TEST_SUITE_P(MpiBackends, TransposeCorrectnessTest, ::testing::ValuesIn(transposeCasesForLabel("mpi")),
                          paramName);
+#if !CUDECOMP_BUILD_CPU_ONLY
 INSTANTIATE_TEST_SUITE_P(NcclBackends, TransposeCorrectnessTest, ::testing::ValuesIn(transposeCasesForLabel("nccl")),
                          paramName);
 INSTANTIATE_TEST_SUITE_P(NvshmemBackends, TransposeCorrectnessTest,
@@ -863,3 +892,5 @@ INSTANTIATE_TEST_SUITE_P(ExternalCudaGraphNvshmemBackends, ExternalCudaGraphTran
                          ::testing::ValuesIn(externalCudaGraphTransposeCases("nvshmem")), paramName);
 INSTANTIATE_TEST_SUITE_P(NcclUserBufferRegistration, NcclUserBufferRegistrationTest,
                          ::testing::ValuesIn(ncclUserBufferRegistrationCases()), paramName);
+
+#endif