espressomd · jngrad · Jul 7, 2026 · Jul 7, 2026
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -209,6 +209,10 @@ set_property(
   TARGET espresso_compiler_flags APPEND
   PROPERTY INTERFACE_COMPILE_FEATURES cxx_std_${ESPRESSO_MINIMAL_CXX_STANDARD})
 
+add_library(espresso_kokkos_compiler_flags INTERFACE)
+add_library(espresso::kokkos::compiler_flags ALIAS
+            espresso_kokkos_compiler_flags)
+
 #
 # AVX2 support
 #
@@ -280,20 +284,35 @@ if(ESPRESSO_BUILD_WITH_CUDA)
   endif()
   if(NOT DEFINED ESPRESSO_CMAKE_CUDA_ARCHITECTURES)
     if("$ENV{CUDAARCHS}" STREQUAL "")
-      # 1. sm_61: GTX-1000 series (Pascal)
-      # 2. sm_75: RTX-2000 series (Turing)
-      # 3. sm_86: RTX-3000 series (Ampere)
-      # 4. sm_89: RTX-4000 series (Ada)
-      # 5. sm_90: H100 series (Hopper)
-      # 6. sm_120: RTX-5000 series (Blackwell)
+      # cmake-format: off
+      # sm_60: P100 series (Pascal)
+      # sm_61: GTX-1000 series (Pascal)
+      # sm_70: V100 series (Volta)
+      # sm_75: RTX-2000 series (Turing)
+      # sm_86: RTX-3000 series (Ampere)
+      # sm_89: RTX-4000 series (Ada)
+      # sm_90: H100 series (Hopper)
+      # sm_120: RTX-5000 series (Blackwell)
+      # cmake-format: on
       set(ESPRESSO_CUDA_ARCHITECTURES "75;86;89")
+      if(Kokkos_ENABLE_CUDA)
+        set(ESPRESSO_CUDA_ARCHITECTURES "86")
+      endif()
     else()
       set(ESPRESSO_CUDA_ARCHITECTURES "$ENV{CUDAARCHS}")
     endif()
     set(ESPRESSO_CMAKE_CUDA_ARCHITECTURES "${ESPRESSO_CUDA_ARCHITECTURES}"
         CACHE INTERNAL "")
   endif()
   set(CMAKE_CUDA_ARCHITECTURES "${ESPRESSO_CMAKE_CUDA_ARCHITECTURES}")
+  list(LENGTH CMAKE_CUDA_ARCHITECTURES ESPRESSO_CMAKE_CUDA_ARCHITECTURES_LEN)
+  if(Kokkos_ENABLE_CUDA AND ESPRESSO_CMAKE_CUDA_ARCHITECTURES_LEN GREATER 1
+     AND NOT CMAKE_CUDA_COMPILER_ID STREQUAL "Clang")
+    message(
+      FATAL_ERROR
+        "Kokkos can only be built against a single CUDA architecture when using the nvcc_wrapper, got CMAKE_CUDA_ARCHITECTURES=${CMAKE_CUDA_ARCHITECTURES}"
+    )
+  endif()
   cmake_path(GET CUDA_cuda_driver_LIBRARY PARENT_PATH ESPRESSO_LIBCUDA_RPATH)
   cmake_path(GET CUDA_cudart_LIBRARY PARENT_PATH ESPRESSO_LIBCUDART_RPATH)
   macro(espresso_add_cuda_rpaths)
@@ -465,8 +484,8 @@ target_compile_options(
     $<$<AND:$<BOOL:${ESPRESSO_BUILD_WITH_HDF5}>,$<COMPILE_LANG_AND_ID:CXX,GNU,Clang,CrayClang,AppleClang,IntelLLVM>>:-Wno-old-style-cast>
     $<$<AND:$<BOOL:${ESPRESSO_BUILD_WITH_HDF5}>,$<COMPILE_LANG_AND_ID:CXX,NVHPC>>:--diag_suppress=storage_class_not_first>
     $<$<AND:$<VERSION_LESS:${ESPRESSO_MINIMAL_CXX_STANDARD},23>,$<COMPILE_LANG_AND_ID:CXX,NVHPC>>:--diag_suppress=bad_pp_directive_keyword>
-    $<$<COMPILE_LANG_AND_ID:CXX,NVHPC>:--diag_suppress=code_is_unreachable>
-    $<$<COMPILE_LANG_AND_ID:CXX,NVHPC>:--diag_suppress=loop_not_reachable>
+    $<$<COMPILE_LANG_AND_ID:CXX,NVHPC,NVIDIA>:--diag_suppress=code_is_unreachable>
+    $<$<COMPILE_LANG_AND_ID:CXX,NVHPC,NVIDIA>:--diag_suppress=loop_not_reachable>
     # disable GCC's C++17 processor-specific ABI warning (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=94383)
     $<$<AND:$<COMPILE_LANG_AND_ID:CXX,GNU>,$<IN_LIST:${CMAKE_SYSTEM_PROCESSOR},arm;arm64;aarch64;aarch64_be;powerpc64le>>:-Wno-psabi>
     # warnings are errors
@@ -701,10 +720,11 @@ if(NOT DEFINED Kokkos_FOUND OR NOT ${Kokkos_FOUND})
     OVERRIDE_FIND_PACKAGE
   )
   # cmake-format: on
-  set(BUILD_SHARED_LIBS ON)
+  set(BUILD_SHARED_LIBS OFF)
   set(CMAKE_SHARED_LIBRARY_PREFIX "lib")
   set(Kokkos_ENABLE_SERIAL ON CACHE BOOL "")
   set(Kokkos_ENABLE_OPENMP ON CACHE BOOL "")
+  set(Kokkos_ENABLE_CUDA_RELOCATABLE_DEVICE_CODE OFF CACHE BOOL "")
   set(Kokkos_ENABLE_IMPL_VIEW_LEGACY ON CACHE BOOL "")
   set(Kokkos_ENABLE_COMPLEX_ALIGN ON CACHE BOOL "")
   set(Kokkos_ENABLE_AGGRESSIVE_VECTORIZATION ON CACHE BOOL "")
@@ -725,6 +745,11 @@ if(NOT DEFINED Kokkos_FOUND OR NOT ${Kokkos_FOUND})
               LIBRARY DESTINATION "${ESPRESSO_INSTALL_LIBDIR}")
     endif()
   endforeach()
+  target_compile_options(
+    kokkoscore
+    PRIVATE
+      $<$<COMPILE_LANG_AND_ID:CXX,NVHPC>:--diag_suppress=initialization_not_reachable>
+  )
   # mark kokkos headers as system headers to disable compiler diagnostics
   set_property(
     TARGET kokkos APPEND
@@ -817,7 +842,7 @@ if(ESPRESSO_BUILD_WITH_NLOPT)
     FetchContent_Declare(
       nlopt
       GIT_REPOSITORY https://github.com/stevengj/nlopt.git
-      GIT_TAG        v2.10.0
+      GIT_TAG        v2.11.0
     )
     # cmake-format: on
     set(BUILD_SHARED_LIBS off)
@@ -831,6 +856,13 @@ if(ESPRESSO_BUILD_WITH_NLOPT)
     set(NLOPT_LUKSAN off CACHE BOOL "")
     FetchContent_MakeAvailable(nlopt)
     set(BUILD_SHARED_LIBS ${ESPRESSO_BUILD_SHARED_LIBS_DEFAULT})
+    target_compile_options(
+      nlopt
+      PRIVATE
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=code_is_unreachable>
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=integer_sign_change>
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=deprecated_entity>
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=mixed_enum_type>)
   endif()
 endif()
 
@@ -842,6 +874,12 @@ if(ESPRESSO_BUILD_WITH_VALGRIND)
   endif()
 endif()
 
+if(Kokkos_ENABLE_CUDA AND NOT CMAKE_CXX_COMPILER_ID STREQUAL Clang)
+  # allow constexpr __host__ function calls in __device__ functions
+  target_compile_options(espresso_kokkos_compiler_flags
+                         INTERFACE -expt-relaxed-constexpr)
+endif()
+
 #
 # MPI
 #
@@ -1038,6 +1076,12 @@ if(ESPRESSO_BUILD_WITH_WALBERLA)
         -Wno-unused-but-set-variable)
     target_link_libraries(walberla_sqlite
                           PRIVATE espresso_walberla_sqlite_compiler_flags)
+    target_compile_options(
+      sqlite3
+      PRIVATE
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=code_is_unreachable>
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=cast_to_qualified_type>
+        $<$<COMPILE_LANG_AND_ID:C,NVHPC>:--diag_suppress=set_but_not_used>)
   endif()
   add_library(espresso_walberla_deps INTERFACE)
   add_library(espresso::walberla_deps ALIAS espresso_walberla_deps)
@@ -1113,7 +1157,7 @@ function(espresso_set_common_target_properties)
     if(ESPRESSO_BUILD_WITH_CUDA)
       if(CMAKE_CUDA_COMPILER_ID STREQUAL "NVIDIA")
         set_target_properties(${TARGET_NAME}
-                              PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+                              PROPERTIES CUDA_SEPARABLE_COMPILATION OFF)
       elseif(CMAKE_CUDA_COMPILER_ID STREQUAL "Clang")
         set_target_properties(${TARGET_NAME} PROPERTIES LINKER_LANGUAGE "CXX")
       endif()

diff --git a/libs/Random123-1.09/include/Random123/features/nvccfeatures.h b/libs/Random123-1.09/include/Random123/features/nvccfeatures.h
@@ -59,7 +59,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 //for the device function
 #ifdef  __CUDA_ARCH__
 #ifndef R123_CUDA_DEVICE
-#define R123_CUDA_DEVICE __device__
+#define R123_CUDA_DEVICE __host__ __device__
 #endif
 
 #ifndef R123_USE_MULHILO64_CUDA_INTRIN

diff --git a/src/core/CMakeLists.txt b/src/core/CMakeLists.txt
@@ -79,7 +79,7 @@ target_link_libraries(
   PUBLIC Cabana::Core Kokkos::kokkos OpenMP::OpenMP_CXX
          $<$<BOOL:${ESPRESSO_BUILD_WITH_CUDA}>:OpenMP::OpenMP_CUDA>
   PRIVATE espresso::config espresso::utils::mpi espresso::shapes
-          espresso::compiler_flags
+          espresso::compiler_flags espresso::kokkos::compiler_flags
           $<$<BOOL:${ESPRESSO_BUILD_WITH_WALBERLA}>:espresso::walberla>
           $<$<BOOL:${ESPRESSO_BUILD_WITH_NLOPT}>:nlopt>
           $<$<BOOL:${ESPRESSO_BUILD_WITH_GSL}>:GSL::gsl>

diff --git a/src/core/MpiCallbacks.hpp b/src/core/MpiCallbacks.hpp
@@ -133,24 +133,27 @@ struct callback_void_t final : public callback_concept_t {
   }
 };
 
-template <class F, class R, class... Args> struct FunctorTypes {
-  using functor_type = F;
-  using return_type = R;
+/** @brief Type traits for a functor. */
+template <class T> struct FunctorTypes;
+
+/** @brief Type traits for an immutable lambda. */
+template <class Class, class Ret, class... Args>
+struct FunctorTypes<Ret (Class::*)(Args...) const> {
+  using functor_type = Class;
+  using return_type = Ret;
   using argument_types = std::tuple<Args...>;
 };
 
-template <class C, class R, class... Args>
-auto functor_types_impl(R (C::*)(Args...) const) {
-  return FunctorTypes<C, R, Args...>{};
-}
+template <class Class, class Ret, class... Args>
+using functor_types_from_args = FunctorTypes<Ret (Class::*)(Args...) const>;
 
 template <class F>
-using functor_types =
-    decltype(functor_types_impl(&std::remove_reference_t<F>::operator()));
+using functor_types_from_lambda =
+    FunctorTypes<decltype(&std::remove_reference_t<F>::operator())>;
 
-template <class CRef, class C, class R, class... Args>
-auto make_model_impl(CRef &&c, FunctorTypes<C, R, Args...>) {
-  return std::make_unique<callback_void_t<C, Args...>>(std::forward<CRef>(c));
+template <class F, class C, class R, class... Args>
+auto make_model_impl(F &&f, functor_types_from_args<C, R, Args...>) {
+  return std::make_unique<callback_void_t<C, Args...>>(std::forward<F>(f));
 }
 
 /**
@@ -160,7 +163,7 @@ auto make_model_impl(CRef &&c, FunctorTypes<C, R, Args...>) {
  * to exist and can not be overloaded.
  */
 template <typename F> auto make_model(F &&f) {
-  return make_model_impl(std::forward<F>(f), functor_types<F>{});
+  return make_model_impl(std::forward<F>(f), functor_types_from_lambda<F>{});
 }
 
 /**
@@ -190,8 +193,9 @@ class MpiCallbacks {
   template <class... Args> class CallbackHandle {
   public:
     template <typename F>
-      requires(std::is_same_v<typename detail::functor_types<F>::argument_types,
-                              std::tuple<Args...>>)
+      requires std::is_same_v<typename detail::functor_types_from_lambda<
+                                  F>::argument_types,
+                              std::tuple<Args...>>
     CallbackHandle(std::shared_ptr<MpiCallbacks> cb, F &&f)
         : m_id(cb->add(std::forward<F>(f))), m_cb(std::move(cb)) {}
 
@@ -216,8 +220,7 @@ class MpiCallbacks {
     auto operator()(ArgRef &&...args) const
         /* Enable if a hypothetical function with signature void(Args..)
          * could be called with the provided arguments. */
-      requires(std::is_void_v<decltype(std::declval<void (*)(Args...)>()(
-                   std::forward<ArgRef>(args)...))>)
+      requires std::is_invocable_r_v<void, void (*)(Args...), ArgRef &&...>
     {
       if (m_cb)
         m_cb->call(m_id, std::forward<ArgRef>(args)...);

diff --git a/src/core/aosoa_pack.hpp b/src/core/aosoa_pack.hpp
@@ -124,11 +124,11 @@ struct CellStructure::AoSoA_pack {
       std::size_t i) const {
     Utils::Vector<T, N> result;
     auto const data = result.data();
-#if !defined(__NVCOMPILER)
-#if (defined(__GNUC__) or defined(__GNUG__)) && !defined(__clang__)
-#pragma GCC unroll 8
-#else
+#if !defined(__NVCOMPILER) && !defined(__CUDACC__)
+#if defined(__clang__)
 #pragma omp unroll
+#elif defined(__GNUC__) or defined(__GNUG__)
+#pragma GCC unroll 8
 #endif
 #endif
     for (std::size_t j = 0ul; j < N; j += 1ul) {
@@ -141,11 +141,11 @@ struct CellStructure::AoSoA_pack {
   void
   set_vector_at(Kokkos::View<T *[N], array_layout, Kokkos::HostSpace> &view,
                 std::size_t i, Utils::Vector<T, N> const &value) {
-#if !defined(__NVCOMPILER)
-#if (defined(__GNUC__) or defined(__GNUG__)) && !defined(__clang__)
-#pragma GCC unroll 8
-#else
+#if !defined(__NVCOMPILER) && !defined(__CUDACC__)
+#if defined(__clang__)
 #pragma omp unroll
+#elif defined(__GNUC__) or defined(__GNUG__)
+#pragma GCC unroll 8
 #endif
 #endif
     for (std::size_t j = 0ul; j < N; j += 1ul) {

diff --git a/src/core/cuda/utils.cuh b/src/core/cuda/utils.cuh
@@ -77,4 +77,4 @@ void cuda_check_errors_exit(const dim3 &block, const dim3 &grid,
   cuda_check_errors_exit(_grid, _block, #_function, __FILE__, __LINE__);
 
 #define KERNELCALL(_function, _grid, _block, ...)                              \
-  KERNELCALL_shared(_function, _grid, _block, 0, ##__VA_ARGS__)
+  KERNELCALL_shared(_function, _grid, _block, 0 __VA_OPT__(, ) __VA_ARGS__)
diff --git a/src/core/p3m/CMakeLists.txt b/src/core/p3m/CMakeLists.txt
@@ -35,6 +35,7 @@ if(ESPRESSO_BUILD_WITH_FFTW)
     espresso_p3m
     PRIVATE espresso::compiler_flags
             espresso::p3m::compiler_flags
+            espresso::kokkos::compiler_flags
             espresso::instrumentation
             espresso::utils
             espresso::config

diff --git a/src/core/p3m/FFTBuffersLegacy.cpp b/src/core/p3m/FFTBuffersLegacy.cpp
@@ -28,11 +28,13 @@
 #include "communication.hpp"
 #include "p3m/common.hpp"
 
+#include <utils/device_qualifier.hpp>
+
 #include <array>
 #include <span>
 
 template <typename FloatType>
-FFTBuffersLegacy<FloatType>::~FFTBuffersLegacy() = default;
+HOST_ONLY_QUALIFIER FFTBuffersLegacy<FloatType>::~FFTBuffersLegacy() = default;
 
 template <typename FloatType>
 void FFTBuffersLegacy<FloatType>::update_mesh_views(

diff --git a/src/core/p3m/FFTBuffersLegacy.hpp b/src/core/p3m/FFTBuffersLegacy.hpp
@@ -31,6 +31,8 @@
 
 #include "fft/vector.hpp"
 
+#include <utils/device_qualifier.hpp>
+
 #include <array>
 #include <type_traits>
 
@@ -49,7 +51,7 @@ class FFTBuffersLegacy : public FFTBuffers<FloatType> {
   std::array<fft::vector<FloatType>, 3u> rs_mesh_fields;
 
 public:
-  ~FFTBuffersLegacy() override;
+  HOST_ONLY_QUALIFIER ~FFTBuffersLegacy() override;
   void init_halo() override;
   void init_meshes(int ca_mesh_size) override;
   void perform_vector_halo_gather() override;

diff --git a/src/core/p3m/data_struct.hpp b/src/core/p3m/data_struct.hpp
@@ -27,6 +27,8 @@
 
 #include "common.hpp"
 
+#include <utils/device_qualifier.hpp>
+
 #include <array>
 #include <cassert>
 #include <memory>
@@ -105,7 +107,7 @@ template <typename FloatType> class FFTBuffers {
 public:
   explicit FFTBuffers(P3MLocalMesh const &local_mesh)
       : local_mesh{local_mesh} {}
-  virtual ~FFTBuffers() = default;
+  virtual HOST_ONLY_QUALIFIER ~FFTBuffers() = default;
   /** @brief Initialize the meshes. */
   virtual void init_meshes(int ca_mesh_size) = 0;
   /** @brief Initialize the halo buffers. */