Add a device-less backend

.github/workflows/test.yml

@@ -25,6 +25,15 @@ jobs:
           - Release
           - Debug
         setup:
+          - arch: none
+            backend: none
+            cc: gcc-13
+            cxx: g++-13
+            fc: gfortran-13
+            container: seissol/gha-cpu:davschneller-gpu-image
+            runner: ubuntu-24.04
+            pythonbreak: true
+            test: true
           - arch: sm_86
             backend: cuda
             cc: gcc-13
@@ -148,3 +157,20 @@ jobs:
           cd build
 
           ./tests
+
+      - id: example-basic
+        name: example-basic
+        if: ${{matrix.setup.backend == 'none'}}
+        run: |
+          cd examples/basic
+          mkdir build && cd build
+
+          export CC=${{matrix.setup.cc}}
+          export CXX=${{matrix.setup.cxx}}
+
+          cmake .. -GNinja \
+            -DDEVICE_BACKEND=${{matrix.setup.backend}} \
+            -DCMAKE_BUILD_TYPE=${{matrix.build_type}}
+
+          ninja
+          ./basic

CMakeLists.txt

@@ -13,17 +13,18 @@ if (NOT DEFINED DEVICE_BACKEND)
     message(FATAL_ERROR "DEVICE_BACKEND variable has not been provided into the submodule")
 else()
     set(FOUND OFF)
-    foreach(VARIANT cuda hip oneapi hipsycl acpp)
+    foreach(VARIANT none cuda hip oneapi hipsycl acpp)
         if (${DEVICE_BACKEND} STREQUAL ${VARIANT})
             set(FOUND ON)
         endif()
     endforeach()
     if (NOT FOUND)
-        message(FATAL_ERROR "DEVICE_BACKEND must be either cuda, hip, opeapi, acpp, or hipsycl. Given: ${DEVICE_BACKEND}")
+        message(FATAL_ERROR "DEVICE_BACKEND must be either none, cuda, hip, oneapi, acpp, or hipsycl. Given: ${DEVICE_BACKEND}")
     endif()
 endif()
 
-if (NOT DEFINED DEVICE_ARCH)
+# the host backend (DEVICE_BACKEND=none) has no device architecture
+if ((NOT ${DEVICE_BACKEND} STREQUAL "none") AND (NOT DEFINED DEVICE_ARCH))
     message(FATAL_ERROR "DEVICE_ARCH is not defined. "
             "Supported for example: sm_60, sm_61, sm_70, sm_71, gfx906, gfx908, dg1, bdw, skl, Gen8, Gen9, Gen11, Gen12LP")
 endif()
@@ -42,7 +43,10 @@ else()
 endif()
 
 # define device library
-if (${DEVICE_BACKEND} STREQUAL "cuda")
+if (${DEVICE_BACKEND} STREQUAL "none")
+    set(BACKEND_FOLDER "host")
+    include(host.cmake)
+elseif (${DEVICE_BACKEND} STREQUAL "cuda")
     set(BACKEND_FOLDER "cuda")
     include(cuda.cmake)
 elseif(${DEVICE_BACKEND} STREQUAL "hip")

README.md

@@ -22,11 +22,12 @@ resolve with `git submodule init` and `git submodule update`
 
 ## Current implementations
 
-Currently, there are three implementations available
+Currently, there are four implementations available
 
 * Nvidia C CUDA
 * AMD HIP
 * SYCL implemented by Intel oneAPI and [AdaptiveCpp](https://github.com/AdaptiveCpp/AdaptiveCpp)
+* the host, for builds without any GPU
 
 ## Setup and Build
 
@@ -40,6 +41,16 @@ for windows, see the belonging batch reference.
 
 If you want to run the examples, follow the instructions in the belonging package.
 
+### Device options for the host
+
+* use `-DDEVICE_BACKEND:STRING=none` to build without any GPU;
+the API is then implemented on the host
+* there is exactly one device, the host itself: device memory is host memory,
+and all work runs synchronously on the calling thread
+* `DEVICE_ARCH` is not needed
+* Complete example call:
+`cmake .. -DDEVICE_BACKEND:STRING=none`
+
 ### Device options for CUDA
 
 * use `-DDEVICE_BACKEND:STRING=cuda` to build the CUDA implementation
@@ -98,7 +109,9 @@ the new folder but implement it regarding the new API
 * Compile and run the basic folder to get feedback if the basic concepts are working
 * Implement examples/jacobi/src/gpu/kernels for your new API
 * compile and run the jacobi benchmark
-* Now switch to the algorithms package and repeat the procedure
+* Now switch to the algorithms package and repeat the procedure;
+instantiate the member templates with the type lists
+from `algorithms/Instantiations.h`, as the other backends do
 * You can now compile and run the examples in the tests/ folder
 
 ## Add another SYCL compiler

UsmAllocator.h

@@ -23,7 +23,7 @@ class UsmAllocator {
   using difference_type = std::ptrdiff_t;
 
   UsmAllocator() noexcept = delete;
-  UsmAllocator(device::DeviceInstance& instance) noexcept : api(instance.api) {}
+  UsmAllocator(device::DeviceInstance& instance) noexcept : api(&instance.api()) {}
 
   UsmAllocator(const UsmAllocator&) noexcept = default;
   UsmAllocator(UsmAllocator&&) noexcept = default;

algorithms/Instantiations.h

@@ -0,0 +1,65 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef SEISSOLDEVICE_ALGORITHMS_INSTANTIATIONS_H_
+#define SEISSOLDEVICE_ALGORITHMS_INSTANTIATIONS_H_
+
+// The types for which every backend provides the member templates of device::Algorithms.
+//
+// Each list applies the macro given as its argument to every entry. A backend instantiates a
+// member template by passing the matching DEVICE_ALGORITHMS_INSTANTIATE_* macro to its list,
+// inside of namespace device; e.g.
+//
+//   DEVICE_ALGORITHMS_ARRAY_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_SCALE_ARRAY)
+//
+// Thus, all backends export the same symbols, and a consumer that links against one of them
+// links against all of them.
+
+// scaleArray, fillArray
+#define DEVICE_ALGORITHMS_ARRAY_TYPES(X) X(float) X(double) X(int) X(unsigned) X(char)
+
+// setToValue
+#define DEVICE_ALGORITHMS_VALUE_TYPES(X) DEVICE_ALGORITHMS_ARRAY_TYPES(X) X(long) X(unsigned long)
+
+// accumulateBatchedData, compareDataWithHost
+#define DEVICE_ALGORITHMS_FLOATING_TYPES(X) X(float) X(double)
+
+// reduceVector; each entry is (accumulator type, vector element type)
+#define DEVICE_ALGORITHMS_REDUCTION_TYPES(X)                                                       \
+  X(int, int)                                                                                      \
+  X(unsigned, unsigned)                                                                            \
+  X(long, int)                                                                                     \
+  X(unsigned long, unsigned)                                                                       \
+  X(long, long)                                                                                    \
+  X(unsigned long, unsigned long)                                                                  \
+  X(long long, int)                                                                                \
+  X(unsigned long long, unsigned)                                                                  \
+  X(long long, long)                                                                               \
+  X(unsigned long long, unsigned long)                                                             \
+  X(long long, long long)                                                                          \
+  X(unsigned long long, unsigned long long)                                                        \
+  X(float, float)                                                                                  \
+  X(double, float)                                                                                 \
+  X(double, double)
+
+#define DEVICE_ALGORITHMS_INSTANTIATE_SCALE_ARRAY(T)                                               \
+  template void Algorithms::scaleArray<T>(T*, T, size_t, void*);
+
+#define DEVICE_ALGORITHMS_INSTANTIATE_FILL_ARRAY(T)                                                \
+  template void Algorithms::fillArray<T>(T*, T, size_t, void*);
+
+#define DEVICE_ALGORITHMS_INSTANTIATE_SET_TO_VALUE(T)                                              \
+  template void Algorithms::setToValue<T>(T**, T, size_t, size_t, void*);
+
+#define DEVICE_ALGORITHMS_INSTANTIATE_ACCUMULATE_BATCHED_DATA(T)                                   \
+  template void Algorithms::accumulateBatchedData<T>(const T**, T**, size_t, size_t, void*);
+
+#define DEVICE_ALGORITHMS_INSTANTIATE_COMPARE_DATA_WITH_HOST(T)                                    \
+  template void Algorithms::compareDataWithHost<T>(const T*, const T*, size_t, const std::string&);
+
+#define DEVICE_ALGORITHMS_INSTANTIATE_REDUCE_VECTOR(AccT, VecT)                                    \
+  template void Algorithms::reduceVector<AccT, VecT>(                                              \
+      AccT*, const VecT*, bool, size_t, ReductionType, void*);
+
+#endif // SEISSOLDEVICE_ALGORITHMS_INSTANTIATIONS_H_

algorithms/cudahip/ArrayManip.cpp

@@ -5,6 +5,7 @@
 #include "AbstractAPI.h"
 #include "Internals.h"
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 
 #include <cassert>
 #include <cstdint>
@@ -28,22 +29,7 @@ void Algorithms::scaleArray(T* devArray, T scalar, const size_t numElements, voi
   kernel_scaleArray<<<grid, block, 0, stream>>>(devArray, scalar, numElements);
   CHECK_ERR;
 }
-template void Algorithms::scaleArray(float* devArray,
-                                     float scalar,
-                                     const size_t numElements,
-                                     void* streamPtr);
-template void Algorithms::scaleArray(double* devArray,
-                                     double scalar,
-                                     const size_t numElements,
-                                     void* streamPtr);
-template void
-    Algorithms::scaleArray(int* devArray, int scalar, const size_t numElements, void* streamPtr);
-template void Algorithms::scaleArray(unsigned* devArray,
-                                     unsigned scalar,
-                                     const size_t numElements,
-                                     void* streamPtr);
-template void
-    Algorithms::scaleArray(char* devArray, char scalar, const size_t numElements, void* streamPtr);
+DEVICE_ALGORITHMS_ARRAY_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_SCALE_ARRAY)
 
 //--------------------------------------------------------------------------------------------------
 template <typename T>
@@ -63,20 +49,7 @@ void Algorithms::fillArray(T* devArray, const T scalar, const size_t numElements
   kernel_fillArray<<<grid, block, 0, stream>>>(devArray, scalar, numElements);
   CHECK_ERR;
 }
-template void
-    Algorithms::fillArray(float* devArray, float scalar, const size_t numElements, void* streamPtr);
-template void Algorithms::fillArray(double* devArray,
-                                    double scalar,
-                                    const size_t numElements,
-                                    void* streamPtr);
-template void
-    Algorithms::fillArray(int* devArray, int scalar, const size_t numElements, void* streamPtr);
-template void Algorithms::fillArray(unsigned* devArray,
-                                    unsigned scalar,
-                                    const size_t numElements,
-                                    void* streamPtr);
-template void
-    Algorithms::fillArray(char* devArray, char scalar, const size_t numElements, void* streamPtr);
+DEVICE_ALGORITHMS_ARRAY_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_FILL_ARRAY)
 
 //--------------------------------------------------------------------------------------------------
 __global__ void kernel_touchMemory(void* ptr, size_t size, bool clean) {

algorithms/cudahip/BatchManip.cpp

@@ -5,6 +5,7 @@
 #include "AbstractAPI.h"
 #include "Internals.h"
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 
 #include <cassert>
 #include <device.h>
@@ -64,17 +65,7 @@ void Algorithms::accumulateBatchedData(
   CHECK_ERR;
 }
 
-template void Algorithms::accumulateBatchedData(const float** baseSrcPtr,
-                                                float** baseDstPtr,
-                                                size_t elementSize,
-                                                size_t numElements,
-                                                void* streamPtr);
-
-template void Algorithms::accumulateBatchedData(const double** baseSrcPtr,
-                                                double** baseDstPtr,
-                                                size_t elementSize,
-                                                size_t numElements,
-                                                void* streamPtr);
+DEVICE_ALGORITHMS_FLOATING_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_ACCUMULATE_BATCHED_DATA)
 
 //--------------------------------------------------------------------------------------------------
 __global__ void
@@ -123,23 +114,7 @@ void Algorithms::setToValue(
   CHECK_ERR;
 }
 
-template void Algorithms::setToValue(
-    float** out, float value, size_t elementSize, size_t numElements, void* streamPtr);
-template void Algorithms::setToValue(
-    double** out, double value, size_t elementSize, size_t numElements, void* streamPtr);
-template void Algorithms::setToValue(
-    int** out, int value, size_t elementSize, size_t numElements, void* streamPtr);
-template void Algorithms::setToValue(
-    unsigned** out, unsigned value, size_t elementSize, size_t numElements, void* streamPtr);
-template void Algorithms::setToValue(
-    long** out, long value, size_t elementSize, size_t numElements, void* streamPtr);
-template void Algorithms::setToValue(unsigned long** out,
-                                     unsigned long value,
-                                     size_t elementSize,
-                                     size_t numElements,
-                                     void* streamPtr);
-template void Algorithms::setToValue(
-    char** out, char value, size_t elementSize, size_t numElements, void* streamPtr);
+DEVICE_ALGORITHMS_VALUE_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_SET_TO_VALUE)
 
 //--------------------------------------------------------------------------------------------------
 __global__ void kernel_copyUniformToScatter(

algorithms/cudahip/Debugging.cpp

@@ -4,6 +4,7 @@
 
 #include "AbstractAPI.h"
 #include "Internals.h"
+#include "algorithms/Instantiations.h"
 #include "utils/logger.h"
 
 #include <cassert>
@@ -46,13 +47,6 @@ void Algorithms::compareDataWithHost(const T* hostPtr,
   delete[] temp;
 };
 
-template void Algorithms::compareDataWithHost(const float* hostPtr,
-                                              const float* devPtr,
-                                              const size_t numElements,
-                                              const std::string& dataName);
-template void Algorithms::compareDataWithHost(const double* hostPtr,
-                                              const double* devPtr,
-                                              const size_t numElements,
-                                              const std::string& dataName);
+DEVICE_ALGORITHMS_FLOATING_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_COMPARE_DATA_WITH_HOST)
 
 } // namespace device