diff --git a/.github/workflows/test.yml b/.github/workflows/test.yml
index bdd29f5..133a8e7 100644
--- a/.github/workflows/test.yml
+++ b/.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
diff --git a/CMakeLists.txt b/CMakeLists.txt
index e58157a..f047e64 100644
--- a/CMakeLists.txt
+++ b/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")
diff --git a/README.md b/README.md
index 7832ef8..69e7f4e 100644
--- a/README.md
+++ b/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
diff --git a/UsmAllocator.h b/UsmAllocator.h
index abb9efb..d971b91 100644
--- a/UsmAllocator.h
+++ b/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;
diff --git a/algorithms/Instantiations.h b/algorithms/Instantiations.h
new file mode 100644
index 0000000..5940f70
--- /dev/null
+++ b/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_
diff --git a/algorithms/cudahip/ArrayManip.cpp b/algorithms/cudahip/ArrayManip.cpp
index 91820ab..f920db6 100644
--- a/algorithms/cudahip/ArrayManip.cpp
+++ b/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) {
diff --git a/algorithms/cudahip/BatchManip.cpp b/algorithms/cudahip/BatchManip.cpp
index 9ce8b04..a393531 100644
--- a/algorithms/cudahip/BatchManip.cpp
+++ b/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(
diff --git a/algorithms/cudahip/Debugging.cpp b/algorithms/cudahip/Debugging.cpp
index d4ce907..b380355 100644
--- a/algorithms/cudahip/Debugging.cpp
+++ b/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
diff --git a/algorithms/cudahip/Reduction.cpp b/algorithms/cudahip/Reduction.cpp
index 12f1bf7..76590fc 100644
--- a/algorithms/cudahip/Reduction.cpp
+++ b/algorithms/cudahip/Reduction.cpp
@@ -4,6 +4,7 @@
 
 #include "Internals.h"
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 
 #include <cassert>
 #include <device.h>
@@ -198,95 +199,6 @@ void Algorithms::reduceVector(AccT* result,
   CHECK_ERR;
 }
 
-template void Algorithms::reduceVector(int* result,
-                                       const int* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned* result,
-                                       const unsigned* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long* result,
-                                       const int* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long* result,
-                                       const unsigned* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long* result,
-                                       const long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long* result,
-                                       const unsigned long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long long* result,
-                                       const int* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long long* result,
-                                       const unsigned* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long long* result,
-                                       const long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long long* result,
-                                       const unsigned long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long long* result,
-                                       const long long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long long* result,
-                                       const unsigned long long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(float* result,
-                                       const float* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(double* result,
-                                       const float* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(double* result,
-                                       const double* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
+DEVICE_ALGORITHMS_REDUCTION_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_REDUCE_VECTOR)
 
 } // namespace device
diff --git a/algorithms/host/ArrayManip.cpp b/algorithms/host/ArrayManip.cpp
new file mode 100644
index 0000000..d9e796b
--- /dev/null
+++ b/algorithms/host/ArrayManip.cpp
@@ -0,0 +1,44 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "algorithms/Instantiations.h"
+
+#include <algorithm>
+#include <cstdint>
+#include <cstring>
+#include <device.h>
+
+namespace device {
+template <typename T>
+void Algorithms::scaleArray(T* devArray, T scalar, const size_t numElements, void* /*streamPtr*/) {
+  for (size_t i = 0; i < numElements; ++i) {
+    devArray[i] = static_cast<T>(devArray[i] * scalar);
+  }
+}
+DEVICE_ALGORITHMS_ARRAY_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_SCALE_ARRAY)
+
+template <typename T>
+void Algorithms::fillArray(T* devArray,
+                           const T scalar,
+                           const size_t numElements,
+                           void* /*streamPtr*/) {
+  std::fill_n(devArray, numElements, scalar);
+}
+DEVICE_ALGORITHMS_ARRAY_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_FILL_ARRAY)
+
+void Algorithms::touchMemoryI(void* ptr, size_t size, bool clean, void* /*streamPtr*/) {
+  // without cleaning, there is nothing to do: host memory needs no migration
+  if (clean) {
+    std::memset(ptr, 0, size);
+  }
+}
+
+void Algorithms::incrementalAddI(
+    void** out, void* base, size_t increment, size_t numElements, void* /*streamPtr*/) {
+  const auto baseAddress = reinterpret_cast<uintptr_t>(base);
+  for (size_t i = 0; i < numElements; ++i) {
+    out[i] = reinterpret_cast<void*>(baseAddress + i * increment);
+  }
+}
+} // namespace device
diff --git a/algorithms/host/BatchManip.cpp b/algorithms/host/BatchManip.cpp
new file mode 100644
index 0000000..753d7f0
--- /dev/null
+++ b/algorithms/host/BatchManip.cpp
@@ -0,0 +1,84 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "algorithms/Instantiations.h"
+
+#include <algorithm>
+#include <cstring>
+#include <device.h>
+
+namespace device {
+void Algorithms::streamBatchedDataI(const void** baseSrcPtr,
+                                    void** baseDstPtr,
+                                    size_t elementSize,
+                                    size_t numElements,
+                                    void* /*streamPtr*/) {
+  for (size_t i = 0; i < numElements; ++i) {
+    if (baseSrcPtr[i] != nullptr && baseDstPtr[i] != nullptr) {
+      std::memmove(baseDstPtr[i], baseSrcPtr[i], elementSize);
+    }
+  }
+}
+
+template <typename T>
+void Algorithms::accumulateBatchedData(const T** baseSrcPtr,
+                                       T** baseDstPtr,
+                                       size_t elementSize,
+                                       size_t numElements,
+                                       void* /*streamPtr*/) {
+  for (size_t i = 0; i < numElements; ++i) {
+    const T* srcElement = baseSrcPtr[i];
+    T* dstElement = baseDstPtr[i];
+    for (size_t j = 0; j < elementSize; ++j) {
+      dstElement[j] += srcElement[j];
+    }
+  }
+}
+DEVICE_ALGORITHMS_FLOATING_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_ACCUMULATE_BATCHED_DATA)
+
+void Algorithms::touchBatchedMemoryI(
+    void** basePtr, size_t elementSize, size_t numElements, bool clean, void* /*streamPtr*/) {
+  // without cleaning, there is nothing to do: host memory needs no migration
+  if (clean) {
+    for (size_t i = 0; i < numElements; ++i) {
+      if (basePtr[i] != nullptr) {
+        std::memset(basePtr[i], 0, elementSize);
+      }
+    }
+  }
+}
+
+template <typename T>
+void Algorithms::setToValue(
+    T** out, T value, size_t elementSize, size_t numElements, void* /*streamPtr*/) {
+  for (size_t i = 0; i < numElements; ++i) {
+    std::fill_n(out[i], elementSize, value);
+  }
+}
+DEVICE_ALGORITHMS_VALUE_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_SET_TO_VALUE)
+
+void Algorithms::copyUniformToScatterI(const void* src,
+                                       void** dst,
+                                       size_t srcOffset,
+                                       size_t copySize,
+                                       size_t numElements,
+                                       void* /*streamPtr*/) {
+  const auto* srcBytes = static_cast<const char*>(src);
+  for (size_t i = 0; i < numElements; ++i) {
+    std::memmove(dst[i], srcBytes + i * srcOffset, copySize);
+  }
+}
+
+void Algorithms::copyScatterToUniformI(const void** src,
+                                       void* dst,
+                                       size_t dstOffset,
+                                       size_t copySize,
+                                       size_t numElements,
+                                       void* /*streamPtr*/) {
+  auto* dstBytes = static_cast<char*>(dst);
+  for (size_t i = 0; i < numElements; ++i) {
+    std::memmove(dstBytes + i * dstOffset, src[i], copySize);
+  }
+}
+} // namespace device
diff --git a/algorithms/host/Debugging.cpp b/algorithms/host/Debugging.cpp
new file mode 100644
index 0000000..51db324
--- /dev/null
+++ b/algorithms/host/Debugging.cpp
@@ -0,0 +1,47 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "algorithms/Instantiations.h"
+#include "utils/logger.h"
+
+#include <cmath>
+#include <device.h>
+#include <sstream>
+#include <string>
+
+namespace device {
+
+template <typename T>
+void Algorithms::compareDataWithHost(const T* hostPtr,
+                                     const T* devPtr,
+                                     const size_t numElements,
+                                     const std::string& dataName) {
+
+  std::stringstream stream;
+  stream << "DEVICE:: comparing array: " << dataName << '\n';
+
+  constexpr T Eps = 1e-12;
+  for (size_t i = 0; i < numElements; ++i) {
+    if (std::abs(hostPtr[i] - devPtr[i]) > Eps) {
+      if ((std::isnan(hostPtr[i])) || (std::isnan(devPtr[i]))) {
+        stream << "DEVICE:: results is NAN. Cannot proceed\n";
+        logError() << stream.str();
+      }
+
+      stream << "DEVICE::ERROR:: host and device arrays are different\n";
+      stream << "DEVICE::ERROR:: "
+             << "host value (" << hostPtr[i] << ") | "
+             << "device value (" << devPtr[i] << ") "
+             << "at index " << i << '\n';
+      stream << "DEVICE::ERROR:: Difference = " << (hostPtr[i] - devPtr[i]) << std::endl;
+      logError() << stream.str();
+    }
+  }
+  stream << "DEVICE:: host and device arrays are the same\n";
+  logInfo() << stream.str();
+}
+
+DEVICE_ALGORITHMS_FLOATING_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_COMPARE_DATA_WITH_HOST)
+
+} // namespace device
diff --git a/algorithms/host/Reduction.cpp b/algorithms/host/Reduction.cpp
new file mode 100644
index 0000000..508bdc5
--- /dev/null
+++ b/algorithms/host/Reduction.cpp
@@ -0,0 +1,60 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "algorithms/Instantiations.h"
+
+#include <device.h>
+#include <limits>
+
+namespace device {
+namespace {
+template <typename AccT, typename VecT, typename OperationT>
+void reduce(AccT* result,
+            const VecT* buffer,
+            bool overrideResult,
+            size_t size,
+            AccT neutral,
+            OperationT operation) {
+  AccT accumulator = overrideResult ? neutral : *result;
+  for (size_t i = 0; i < size; ++i) {
+    accumulator = operation(accumulator, static_cast<AccT>(buffer[i]));
+  }
+  *result = accumulator;
+}
+} // namespace
+
+template <typename AccT, typename VecT>
+void Algorithms::reduceVector(AccT* result,
+                              const VecT* buffer,
+                              bool overrideResult,
+                              size_t size,
+                              ReductionType type,
+                              void* /*streamPtr*/) {
+  switch (type) {
+  case ReductionType::Add: {
+    reduce(result, buffer, overrideResult, size, AccT{0}, [](AccT a, AccT b) { return a + b; });
+    break;
+  }
+  case ReductionType::Max: {
+    reduce(result,
+           buffer,
+           overrideResult,
+           size,
+           std::numeric_limits<AccT>::lowest(),
+           [](AccT a, AccT b) { return a > b ? a : b; });
+    break;
+  }
+  case ReductionType::Min: {
+    reduce(
+        result, buffer, overrideResult, size, std::numeric_limits<AccT>::max(), [](AccT a, AccT b) {
+          return a > b ? b : a;
+        });
+    break;
+  }
+  }
+}
+
+DEVICE_ALGORITHMS_REDUCTION_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_REDUCE_VECTOR)
+
+} // namespace device
diff --git a/algorithms/sycl/ArrayManip.cpp b/algorithms/sycl/ArrayManip.cpp
index 4a72d23..eb41025 100644
--- a/algorithms/sycl/ArrayManip.cpp
+++ b/algorithms/sycl/ArrayManip.cpp
@@ -4,6 +4,7 @@
 
 #include "AbstractAPI.h"
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 #include "interfaces/sycl/Internals.h"
 
 #include <device.h>
@@ -25,22 +26,7 @@ void Algorithms::scaleArray(T* devArray, T scalar, size_t numElements, void* str
     });
   });
 }
-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>
 void Algorithms::fillArray(T* devArray, const T scalar, const size_t numElements, void* streamPtr) {
@@ -56,20 +42,7 @@ void Algorithms::fillArray(T* devArray, const T scalar, const size_t numElements
   });
 }
 
-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)
 
 void Algorithms::touchMemoryI(void* ptr, size_t size, bool clean, void* streamPtr) {
   auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, size);
diff --git a/algorithms/sycl/BatchManip.cpp b/algorithms/sycl/BatchManip.cpp
index 0aab2b3..945a6d6 100644
--- a/algorithms/sycl/BatchManip.cpp
+++ b/algorithms/sycl/BatchManip.cpp
@@ -4,6 +4,7 @@
 
 #include "AbstractAPI.h"
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 #include "interfaces/sycl/Internals.h"
 
 #include <device.h>
@@ -55,17 +56,7 @@ void Algorithms::accumulateBatchedData(
   });
 }
 
-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)
 
 void Algorithms::touchBatchedMemoryI(
     void** basePtr, size_t elementSize, size_t numElements, bool clean, void* streamPtr) {
@@ -110,23 +101,7 @@ void Algorithms::setToValue(
   });
 }
 
-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)
 
 void Algorithms::copyUniformToScatterI(const void* src,
                                        void** dst,
diff --git a/algorithms/sycl/Debugging.cpp b/algorithms/sycl/Debugging.cpp
index 9342cf2..0af185c 100644
--- a/algorithms/sycl/Debugging.cpp
+++ b/algorithms/sycl/Debugging.cpp
@@ -4,6 +4,7 @@
 
 #include "AbstractAPI.h"
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 #include "interfaces/sycl/Internals.h"
 #include "utils/logger.h"
 
@@ -47,13 +48,6 @@ void Algorithms::compareDataWithHost(const T* hostPtr,
   logInfo() << stream.str();
   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
diff --git a/algorithms/sycl/Reduction.cpp b/algorithms/sycl/Reduction.cpp
index e0d5d52..df93012 100644
--- a/algorithms/sycl/Reduction.cpp
+++ b/algorithms/sycl/Reduction.cpp
@@ -3,6 +3,7 @@
 // SPDX-License-Identifier: BSD-3-Clause
 
 #include "algorithms/Common.h"
+#include "algorithms/Instantiations.h"
 #include "interfaces/sycl/Internals.h"
 #include "utils/logger.h"
 
@@ -200,95 +201,6 @@ void Algorithms::reduceVector(AccT* result,
 
 #endif
 
-template void Algorithms::reduceVector(int* result,
-                                       const int* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned* result,
-                                       const unsigned* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long* result,
-                                       const int* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long* result,
-                                       const unsigned* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long* result,
-                                       const long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long* result,
-                                       const unsigned long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long long* result,
-                                       const int* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long long* result,
-                                       const unsigned* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long long* result,
-                                       const long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long long* result,
-                                       const unsigned long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(long long* result,
-                                       const long long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(unsigned long long* result,
-                                       const unsigned long long* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(float* result,
-                                       const float* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(double* result,
-                                       const float* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
-template void Algorithms::reduceVector(double* result,
-                                       const double* buffer,
-                                       bool overrideResult,
-                                       size_t size,
-                                       ReductionType type,
-                                       void* streamPtr);
+DEVICE_ALGORITHMS_REDUCTION_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_REDUCE_VECTOR)
 
 } // namespace device
diff --git a/device.cpp b/device.cpp
index f1b5182..51851a6 100644
--- a/device.cpp
+++ b/device.cpp
@@ -10,24 +10,29 @@
 #include "interfaces/hip/HipWrappedAPI.h"
 #elif DEVICE_LANG_SYCL
 #include "interfaces/sycl/SyclWrappedAPI.h"
+#elif DEVICE_LANG_HOST
+#include "interfaces/host/HostWrappedAPI.h"
 #else
 #error "Unknown interface for the device wrapper"
 #endif
 
+#include <memory>
+
 using namespace device;
 
-DeviceInstance::DeviceInstance() {
-  // NOTE: all headers inside of macros define their unique ConcreteInterface.
-  // Make sure to not include multiple different interfaces at the same time.
-  // Only one interface is allowed per program because of issues of unique compilers, etc.
-  api = new ConcreteAPI;
-  algorithms.setDeviceApi(api);
+// NOTE: all headers inside of macros define their unique ConcreteInterface.
+// Make sure to not include multiple different interfaces at the same time.
+// Only one interface is allowed per program because of issues of unique compilers, etc.
+DeviceInstance::DeviceInstance()
+    : apiP(std::make_unique<ConcreteAPI>()), algorithmsP(std::make_unique<Algorithms>()) {
+  algorithmsP->setDeviceApi(apiP.get());
 }
 
-DeviceInstance::~DeviceInstance() {
-  this->finalize();
-  delete api;
-  api = nullptr;
-}
+DeviceInstance::~DeviceInstance() { this->finalize(); }
 
-void DeviceInstance::finalize() { api->finalize(); }
+void DeviceInstance::finalize() { apiP->finalize(); }
+
+DeviceInstance& DeviceInstance::instance() {
+  static DeviceInstance singleton;
+  return singleton;
+}
diff --git a/device.h b/device.h
index e7172b9..b80d34a 100644
--- a/device.h
+++ b/device.h
@@ -8,6 +8,8 @@
 #include "AbstractAPI.h"
 #include "Algorithms.h"
 
+#include <memory>
+
 namespace device {
 
 // DeviceInstance -> Singleton
@@ -15,18 +17,24 @@ class DeviceInstance {
   public:
   DeviceInstance(const DeviceInstance&) = delete;
   DeviceInstance& operator=(const DeviceInstance&) = delete;
-  static DeviceInstance& getInstance() {
-    static DeviceInstance instance;
-    return instance;
-  }
+
+  // defined in device.cpp: all shared objects of a process thus see the same instance, also those
+  // which are compiled with hidden visibility (e.g. Python modules)
+  static DeviceInstance& instance();
+
   ~DeviceInstance();
   void finalize();
 
-  AbstractAPI* api{nullptr};
-  Algorithms algorithms{};
+  // the API and the algorithms exist for the whole lifetime of the instance; the constness of
+  // the instance is shallow, i.e. a const instance gives mutable access to them
+  [[nodiscard]] AbstractAPI& api() const { return *apiP; }
+  [[nodiscard]] Algorithms& algorithms() const { return *algorithmsP; }
 
   private:
   DeviceInstance();
+
+  std::unique_ptr<AbstractAPI> apiP;
+  std::unique_ptr<Algorithms> algorithmsP;
 };
 } // namespace device
 
diff --git a/examples/basic/basic.cpp b/examples/basic/basic.cpp
index ac9d5ea..11bdf49 100644
--- a/examples/basic/basic.cpp
+++ b/examples/basic/basic.cpp
@@ -2,11 +2,11 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
-#include "DataTypes.h"
+#include "common.h"
 #include "device.h"
 
+#include <cstdlib>
 #include <iostream>
-#include <stdlib.h>
 
 using namespace device;
 
@@ -14,48 +14,60 @@ int main(int argc, char* argv[]) {
   const size_t size = 1024;
   real* inputArray = new real[size];
   real* outputArray = new real[size];
+  for (size_t i = 0; i < size; ++i) {
+    inputArray[i] = static_cast<real>(i);
+    outputArray[i] = 0;
+  }
 
-  DeviceInstance& device = DeviceInstance::getInstance();
+  DeviceInstance& device = DeviceInstance::instance();
 
   // set up the first device
-  const int numDevices = device.api->getNumDevices();
+  const int numDevices = device.api().getNumDevices();
   std::cout << "Num. devices available: " << numDevices << '\n';
   if (numDevices > 0) {
-    device.api->setDevice(0);
+    device.api().setDevice(0);
   }
-  device.api->initialize();
+  device.api().initialize();
 
   // print some device info
-  std::string deviceInfo(device.api->getDeviceInfoAsText(0));
+  std::string deviceInfo(device.api().getDeviceInfoAsText(0));
   std::cout << deviceInfo << std::endl;
 
-  std::cout << "alignment: " << device.api->getGlobMemAlignment() << std::endl;
-  std::cout << "max available mem: " << device.api->getMaxAvailableMem() << std::endl;
-  std::cout << "max shared mem: " << device.api->getMaxSharedMemSize() << std::endl;
-  std::cout << "max thread block size: " << device.api->getMaxThreadBlockSize() << std::endl;
+  std::cout << "alignment: " << device.api().getGlobMemAlignment() << std::endl;
+  std::cout << "max available mem: " << device.api().getMaxAvailableMem() << std::endl;
 
   // allocate mem. on a device
-  real* dInputArray = static_cast<real*>(device.api->allocGlobMem(sizeof(real) * size));
-  real* dOutputArray = static_cast<real*>(device.api->allocGlobMem(sizeof(real) * size));
+  real* dInputArray = static_cast<real*>(device.api().allocGlobMem(sizeof(real) * size));
+  real* dOutputArray = static_cast<real*>(device.api().allocGlobMem(sizeof(real) * size));
 
   // copy data into a device
-  device.api->copyTo(dInputArray, inputArray, sizeof(real) * size);
+  device.api().copyTo(dInputArray, inputArray, sizeof(real) * size);
 
-  // call a kernel
-  device.api->checkOffloading();
-  device.api->syncDevice();
+  // copy data on the device
+  device.api().copyBetween(dOutputArray, dInputArray, sizeof(real) * size);
+  device.api().syncDevice();
 
   // copy data from a device
-  device.api->copyFrom(outputArray, dOutputArray, sizeof(real) * size);
+  device.api().copyFrom(outputArray, dOutputArray, sizeof(real) * size);
+
+  size_t mismatches = 0;
+  for (size_t i = 0; i < size; ++i) {
+    if (outputArray[i] != inputArray[i]) {
+      ++mismatches;
+    }
+  }
+  std::cout << "mismatches after the round trip: " << mismatches << std::endl;
 
   // deallocate mem. on a device
-  device.api->freeGlobMem(dInputArray);
-  device.api->freeGlobMem(dOutputArray);
+  device.api().freeGlobMem(dInputArray);
+  device.api().freeGlobMem(dOutputArray);
 
-  std::cout << device.api->getMemLeaksReport();
+  std::cout << device.api().getMemLeaksReport();
 
   device.finalize();
 
   delete[] outputArray;
   delete[] inputArray;
+
+  return mismatches == 0 ? EXIT_SUCCESS : EXIT_FAILURE;
 }
diff --git a/examples/cmake/common_options.cmake b/examples/cmake/common_options.cmake
index 4051204..323604c 100644
--- a/examples/cmake/common_options.cmake
+++ b/examples/cmake/common_options.cmake
@@ -13,7 +13,7 @@ set(REAL_SIZE_IN_BYTES "8" CACHE STRING "size of the floating point data type")
 set_property(CACHE REAL_SIZE_IN_BYTES PROPERTY STRINGS "8" "4")
 
 set(DEVICE_BACKEND "cuda" CACHE STRING "type of an interface")
-set_property(CACHE DEVICE_BACKEND PROPERTY STRINGS "cuda" "hip" "oneapi" "hipsycl")
+set_property(CACHE DEVICE_BACKEND PROPERTY STRINGS "none" "cuda" "hip" "oneapi" "hipsycl" "acpp")
 
 set(DEVICE_ARCH "sm_60" CACHE STRING "model of Streaming Multiprocessor")
 set(DEVICE_ARCH_OPTIONS "sm_50" "sm_60" "sm_70" "sm_71" "sm_75" "sm_80" "sm_86"
diff --git a/examples/jacobi/src/gpu/solver.cpp b/examples/jacobi/src/gpu/solver.cpp
index e7f6fda..69a0127 100644
--- a/examples/jacobi/src/gpu/solver.cpp
+++ b/examples/jacobi/src/gpu/solver.cpp
@@ -41,22 +41,22 @@ void gpu::Solver::run(const SolverSettingsT& settings,
   std::tie(invDiag, lu) = getDLU(matrix);
 
   // allocate gpu data structures
-  DeviceInstance& device = DeviceInstance::getInstance();
-  device.api->setDevice(ws.rank);
-  device.api->initialize();
+  DeviceInstance& device = DeviceInstance::instance();
+  device.api().setDevice(ws.rank);
+  device.api().initialize();
   this->setUp(lu, rhs, x, residual, invDiag);
 
   // start solver
   Statistics computeStat(ws, range);
   Statistics commStat(ws, range);
-  auto defaultStream = device.api->getDefaultStream();
+  auto defaultStream = device.api().getDefaultStream();
   while ((infNorm > settings.eps) and (currentIter <= settings.maxNumIters)) {
 
     computeStat.start();
     launch_multMatVec(*devLU, devX, devTemp, defaultStream);
     launch_manipVectors(range, devRhs, devTemp, devX, VectorManipOps::Subtraction, defaultStream);
     launch_manipVectors(range, devInvDiag, devX, devX, VectorManipOps::Multiply, defaultStream);
-    device.api->syncDevice();
+    device.api().syncDevice();
     computeStat.stop();
 
     commStat.start();
@@ -73,7 +73,7 @@ void gpu::Solver::run(const SolverSettingsT& settings,
           range, devTemp, devResidual, devResidual, VectorManipOps::Addition, defaultStream);
       launch_manipVectors(
           range, devRhs, devResidual, devResidual, VectorManipOps::Subtraction, defaultStream);
-      device.api->copyFrom(
+      device.api().copyFrom(
           const_cast<real*>(residual.data()), devResidual, residual.size() * sizeof(real));
 
       auto localInfNorm = infNorm = host::getInfNorm(range, residual);
@@ -96,9 +96,9 @@ void gpu::Solver::run(const SolverSettingsT& settings,
     ++currentIter;
   }
 
-  device.api->syncDevice();
+  device.api().syncDevice();
   assembler.assemble<SystemType::OnDevice>(devX, devTempX);
-  device.api->copyFrom(const_cast<real*>(x.data()), devTempX, x.size() * sizeof(real));
+  device.api().copyFrom(const_cast<real*>(x.data()), devTempX, x.size() * sizeof(real));
 
   this->tearDown();
 }
@@ -108,22 +108,22 @@ void gpu::Solver::setUp(const CpuMatrixDataT& lu,
                         const VectorT& x,
                         const VectorT& residual,
                         VectorT& invDiag) {
-  DeviceInstance& device = DeviceInstance::getInstance();
+  DeviceInstance& device = DeviceInstance::instance();
 
-  devRhs = static_cast<real*>(device.api->allocGlobMem(rhs.size() * sizeof(real)));
-  device.api->copyTo(devRhs, rhs.data(), rhs.size() * sizeof(real));
+  devRhs = static_cast<real*>(device.api().allocGlobMem(rhs.size() * sizeof(real)));
+  device.api().copyTo(devRhs, rhs.data(), rhs.size() * sizeof(real));
 
-  devX = static_cast<real*>(device.api->allocGlobMem(x.size() * sizeof(real)));
-  device.api->copyTo(devX, x.data(), x.size() * sizeof(real));
+  devX = static_cast<real*>(device.api().allocGlobMem(x.size() * sizeof(real)));
+  device.api().copyTo(devX, x.data(), x.size() * sizeof(real));
 
-  devTempX = static_cast<real*>(device.api->allocGlobMem(x.size() * sizeof(real)));
-  device.api->copyTo(devTempX, x.data(), x.size() * sizeof(real));
+  devTempX = static_cast<real*>(device.api().allocGlobMem(x.size() * sizeof(real)));
+  device.api().copyTo(devTempX, x.data(), x.size() * sizeof(real));
 
-  devTemp = static_cast<real*>(device.api->allocGlobMem(x.size() * sizeof(real)));
+  devTemp = static_cast<real*>(device.api().allocGlobMem(x.size() * sizeof(real)));
 
   // InvDiag still holds the diagonal elements at this point
-  devDiag = static_cast<real*>(device.api->allocGlobMem(invDiag.size() * sizeof(real)));
-  device.api->copyTo(devDiag, invDiag.data(), invDiag.size() * sizeof(real));
+  devDiag = static_cast<real*>(device.api().allocGlobMem(invDiag.size() * sizeof(real)));
+  device.api().copyTo(devDiag, invDiag.data(), invDiag.size() * sizeof(real));
 
   // compute inverse diagonal matrix
   std::transform(invDiag.begin(), invDiag.end(), invDiag.begin(), [](const real& diag) {
@@ -131,32 +131,32 @@ void gpu::Solver::setUp(const CpuMatrixDataT& lu,
     return 1.0 / diag;
   });
 
-  devInvDiag = static_cast<real*>(device.api->allocGlobMem(invDiag.size() * sizeof(real)));
-  device.api->copyTo(devInvDiag, invDiag.data(), invDiag.size() * sizeof(real));
+  devInvDiag = static_cast<real*>(device.api().allocGlobMem(invDiag.size() * sizeof(real)));
+  device.api().copyTo(devInvDiag, invDiag.data(), invDiag.size() * sizeof(real));
 
-  devResidual = static_cast<real*>(device.api->allocGlobMem(residual.size() * sizeof(real)));
-  device.api->copyTo(devResidual, residual.data(), residual.size() * sizeof(real));
+  devResidual = static_cast<real*>(device.api().allocGlobMem(residual.size() * sizeof(real)));
+  device.api().copyTo(devResidual, residual.data(), residual.size() * sizeof(real));
 
   devLU = std::make_unique<GpuMatrixDataT>(lu.info);
   // GpuMatrixDataT devLU(LU.Info);
-  devLU->data = static_cast<real*>(device.api->allocGlobMem(lu.info.volume * sizeof(real)));
-  device.api->copyTo(devLU->data, lu.data.data(), lu.info.volume * sizeof(real));
+  devLU->data = static_cast<real*>(device.api().allocGlobMem(lu.info.volume * sizeof(real)));
+  device.api().copyTo(devLU->data, lu.data.data(), lu.info.volume * sizeof(real));
 
-  devLU->indices = static_cast<int*>(device.api->allocGlobMem(lu.info.volume * sizeof(int)));
-  device.api->copyTo(devLU->indices, lu.indices.data(), lu.info.volume * sizeof(int));
+  devLU->indices = static_cast<int*>(device.api().allocGlobMem(lu.info.volume * sizeof(int)));
+  device.api().copyTo(devLU->indices, lu.indices.data(), lu.info.volume * sizeof(int));
 }
 
 void gpu::Solver::tearDown() {
-  DeviceInstance& device = DeviceInstance::getInstance();
-  device.api->freeGlobMem(devResidual);
-  device.api->freeGlobMem(devInvDiag);
-  device.api->freeGlobMem(devDiag);
-  device.api->freeGlobMem(devTemp);
-  device.api->freeGlobMem(devX);
-  device.api->freeGlobMem(devTempX);
-  device.api->freeGlobMem(devRhs);
-
-  device.api->freeGlobMem(devLU->data);
-  device.api->freeGlobMem(devLU->indices);
+  DeviceInstance& device = DeviceInstance::instance();
+  device.api().freeGlobMem(devResidual);
+  device.api().freeGlobMem(devInvDiag);
+  device.api().freeGlobMem(devDiag);
+  device.api().freeGlobMem(devTemp);
+  device.api().freeGlobMem(devX);
+  device.api().freeGlobMem(devTempX);
+  device.api().freeGlobMem(devRhs);
+
+  device.api().freeGlobMem(devLU->data);
+  device.api().freeGlobMem(devLU->indices);
   devLU.reset(nullptr);
 }
diff --git a/examples/jacobi/src/helper.hpp b/examples/jacobi/src/helper.hpp
index b924d30..a0e8060 100644
--- a/examples/jacobi/src/helper.hpp
+++ b/examples/jacobi/src/helper.hpp
@@ -106,8 +106,8 @@ class VectorAssembler {
                     recvCounts[0] * sizeof(real));
       }
     } else {
-      device::DeviceInstance& device = device::DeviceInstance::getInstance();
-      device.api->copyBetween(dest, src, recvCounts[0] * sizeof(real));
+      device::DeviceInstance& device = device::DeviceInstance::instance();
+      device.api().copyBetween(dest, src, recvCounts[0] * sizeof(real));
     }
 #endif
   }
diff --git a/examples/jacobi/tests/driver.cpp b/examples/jacobi/tests/driver.cpp
index 07fa941..8615dc5 100644
--- a/examples/jacobi/tests/driver.cpp
+++ b/examples/jacobi/tests/driver.cpp
@@ -25,9 +25,9 @@ int main(int argc, char** argv) {
   WorkSpaceT ws{MPI_COMM_WORLD};
 #endif
 
-  DeviceInstance& device = DeviceInstance::getInstance();
-  device.api->setDevice(ws.rank);
-  device.api->initialize();
+  DeviceInstance& device = DeviceInstance::instance();
+  device.api().setDevice(ws.rank);
+  device.api().initialize();
 
   ::testing::InitGoogleTest(&argc, argv);
   int result = RUN_ALL_TESTS();
diff --git a/examples/jacobi/tests/gpu/subroutines_tests.cpp b/examples/jacobi/tests/gpu/subroutines_tests.cpp
index 42e9f4b..3569c3f 100644
--- a/examples/jacobi/tests/gpu/subroutines_tests.cpp
+++ b/examples/jacobi/tests/gpu/subroutines_tests.cpp
@@ -13,7 +13,7 @@
 using ::testing::ElementsAreArray;
 
 TEST(Subroutines, MultMatrixVec) {
-  auto* api = ::device::DeviceInstance::getInstance().api;
+  auto* api = &::device::DeviceInstance::instance().api();
   auto defaultStream = api->getDefaultStream();
   const int size = 3;
 
@@ -54,7 +54,7 @@ TEST(Subroutines, MultMatrixVec) {
 }
 
 TEST(Subroutines, VectorManips) {
-  auto* api = ::device::DeviceInstance::getInstance().api;
+  auto* api = &::device::DeviceInstance::instance().api();
   auto defaultStream = api->getDefaultStream();
 
   const int size = 3;
diff --git a/examples/mpi/src/main.cpp b/examples/mpi/src/main.cpp
index 9700959..42e785c 100644
--- a/examples/mpi/src/main.cpp
+++ b/examples/mpi/src/main.cpp
@@ -10,8 +10,8 @@ using namespace device;
 
 void forkOther(int otherRank) {
 
-  DeviceInstance& device = DeviceInstance::getInstance();
-  auto* api = device.api;
+  DeviceInstance& device = DeviceInstance::instance();
+  auto* api = &device.api();
   api->setDevice(otherRank);
 
   auto* devPtr = (int*)api->allocGlobMem(sizeof(int));
@@ -24,8 +24,8 @@ void forkOther(int otherRank) {
 }
 
 void forkRoot(int rootRank) {
-  DeviceInstance& device = DeviceInstance::getInstance();
-  auto* api = device.api;
+  DeviceInstance& device = DeviceInstance::instance();
+  auto* api = &device.api();
   api->setDevice(rootRank);
 
   auto* devPtr = (int*)api->allocGlobMem(sizeof(int));
diff --git a/host.cmake b/host.cmake
new file mode 100644
index 0000000..f9ab5bf
--- /dev/null
+++ b/host.cmake
@@ -0,0 +1,14 @@
+# SPDX-FileCopyrightText: 2026 SeisSol Group
+#
+# SPDX-License-Identifier: BSD-3-Clause
+
+add_library(device ${DEVICE_LIBTYPE} device.cpp
+                 interfaces/host/Control.cpp
+                 interfaces/host/Memory.cpp
+                 interfaces/host/Streams.cpp
+                 algorithms/host/ArrayManip.cpp
+                 algorithms/host/BatchManip.cpp
+                 algorithms/host/Debugging.cpp
+                 algorithms/host/Reduction.cpp)
+
+set_target_properties(device PROPERTIES POSITION_INDEPENDENT_CODE ON)
diff --git a/interfaces/host/Control.cpp b/interfaces/host/Control.cpp
new file mode 100644
index 0000000..0b432a3
--- /dev/null
+++ b/interfaces/host/Control.cpp
@@ -0,0 +1,107 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "HostWrappedAPI.h"
+#include "utils/logger.h"
+
+#include <sstream>
+#include <string>
+#include <thread>
+#include <unistd.h>
+
+using namespace device;
+
+namespace {
+// the same alignment as the GPU backends use; it keeps data layouts identical across backends
+constexpr unsigned Alignment = 128;
+
+void checkDeviceId(int deviceId) {
+  if (deviceId != 0) {
+    logError() << "The host backend has exactly one device (id 0); requested was" << deviceId;
+  }
+}
+} // namespace
+
+void ConcreteAPI::setDevice(int deviceId) {
+  checkDeviceId(deviceId);
+  status[StatusID::DriverApiInitialized] = true;
+  status[StatusID::DeviceSelected] = true;
+}
+
+int ConcreteAPI::getDeviceId() {
+  if (!status[StatusID::DeviceSelected]) {
+    logError() << "Device has not been selected. Please, select device before requesting device Id";
+  }
+  return 0;
+}
+
+int ConcreteAPI::getNumDevices() { return 1; }
+
+unsigned ConcreteAPI::getGlobMemAlignment() { return Alignment; }
+
+std::string ConcreteAPI::getDeviceInfoAsText(int deviceId) {
+  checkDeviceId(deviceId);
+  std::ostringstream info;
+  info << "name: " << getDeviceName(deviceId) << '\n';
+  info << "execution: synchronous, on the calling host thread\n";
+  info << "hardwareConcurrency: " << std::thread::hardware_concurrency() << '\n';
+  info << "totalGlobalMem: " << getMaxAvailableMem() << '\n';
+  info << "globMemAlignment: " << getGlobMemAlignment() << '\n';
+  return info.str();
+}
+
+void ConcreteAPI::syncDevice() { isFlagSet<DeviceSelected>(status); }
+
+std::string ConcreteAPI::getApiName() { return "Host"; }
+
+std::string ConcreteAPI::getDeviceName(int deviceId) {
+  checkDeviceId(deviceId);
+  return "Host";
+}
+
+std::string ConcreteAPI::getPciAddress(int deviceId) {
+  checkDeviceId(deviceId);
+  return "";
+}
+
+size_t ConcreteAPI::getMaxAvailableMem() {
+  const long pages = sysconf(_SC_PHYS_PAGES);
+  const long pageSize = sysconf(_SC_PAGE_SIZE);
+  if (pages <= 0 || pageSize <= 0) {
+    return 0;
+  }
+  return static_cast<size_t>(pages) * static_cast<size_t>(pageSize);
+}
+
+bool ConcreteAPI::isUnifiedMemoryDefault() {
+  // the host accesses all memory of this backend directly
+  return true;
+}
+
+void ConcreteAPI::initialize() {
+  if (!status[StatusID::DeviceSelected]) {
+    logError() << "Device has not been selected. Please, select device before calling initialize";
+  }
+  status[StatusID::InterfaceInitialized] = true;
+}
+
+void ConcreteAPI::finalize() {
+  if (status[StatusID::InterfaceInitialized]) {
+    if (!genericStreams.empty()) {
+      logInfo() << "DEVICE::WARNING:" << genericStreams.size()
+                << "device generic stream(s) were not deleted.";
+      genericStreams.clear();
+    }
+    status[StatusID::InterfaceInitialized] = false;
+  }
+  m_isFinalized = true;
+}
+
+void ConcreteAPI::profilingMessage(const std::string& /*message*/) {}
+
+void ConcreteAPI::putProfilingMark(const std::string& /*name*/, ProfilingColors /*color*/) {}
+
+void ConcreteAPI::popLastProfilingMark() {}
+
+void ConcreteAPI::setupPrinting(int /*rank*/) {}
diff --git a/interfaces/host/HostWrappedAPI.h b/interfaces/host/HostWrappedAPI.h
new file mode 100644
index 0000000..dc25e42
--- /dev/null
+++ b/interfaces/host/HostWrappedAPI.h
@@ -0,0 +1,147 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef SEISSOLDEVICE_INTERFACES_HOST_HOSTWRAPPEDAPI_H_
+#define SEISSOLDEVICE_INTERFACES_HOST_HOSTWRAPPEDAPI_H_
+
+#include "AbstractAPI.h"
+#include "Common.h"
+
+#include <chrono>
+#include <cstddef>
+#include <cstdint>
+#include <functional>
+#include <memory>
+#include <string>
+#include <unordered_map>
+#include <vector>
+
+namespace device {
+
+/**
+ * The device API, 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; that is, every operation has completed once its call
+ * returns. Streams and events only serve as handles; waiting on them returns immediately.
+ */
+class ConcreteAPI : public AbstractAPI {
+  public:
+  void setDevice(int deviceId) override;
+  int getDeviceId() override;
+
+  int getNumDevices() override;
+  unsigned getGlobMemAlignment() override;
+  std::string getDeviceInfoAsText(int deviceId) override;
+  void syncDevice() override;
+
+  std::string getApiName() override;
+  std::string getDeviceName(int deviceId) override;
+  std::string getPciAddress(int deviceId) override;
+
+  void* allocGlobMem(size_t size, bool compress) override;
+  void* allocUnifiedMem(size_t size, bool compress, Destination hint) override;
+  void* allocPinnedMem(size_t size, bool compress, Destination hint) override;
+  void freeGlobMem(void* devPtr) override;
+  void freeUnifiedMem(void* devPtr) override;
+  void freePinnedMem(void* devPtr) override;
+  std::string getMemLeaksReport() override;
+
+  void* allocMemAsync(size_t size, void* streamPtr) override;
+  void freeMemAsync(void* devPtr, void* streamPtr) override;
+
+  void pinMemory(void* ptr, size_t size) override;
+  void unpinMemory(void* ptr) override;
+  void* devicePointer(void* ptr) override;
+
+  void copyTo(void* dst, const void* src, size_t count) override;
+  void copyFrom(void* dst, const void* src, size_t count) override;
+  void copyBetween(void* dst, const void* src, size_t count) override;
+  void copyToAsync(void* dst, const void* src, size_t count, void* streamPtr) override;
+  void copyFromAsync(void* dst, const void* src, size_t count, void* streamPtr) override;
+  void copyBetweenAsync(void* dst, const void* src, size_t count, void* streamPtr) override;
+
+  void copy2dArrayTo(void* dst,
+                     size_t dpitch,
+                     const void* src,
+                     size_t spitch,
+                     size_t width,
+                     size_t height) override;
+  void copy2dArrayFrom(void* dst,
+                       size_t dpitch,
+                       const void* src,
+                       size_t spitch,
+                       size_t width,
+                       size_t height) override;
+  void prefetchUnifiedMemTo(Destination type,
+                            const void* devPtr,
+                            size_t count,
+                            void* streamPtr) override;
+
+  size_t getMaxAvailableMem() override;
+  size_t getCurrentlyOccupiedMem() override;
+  size_t getCurrentlyOccupiedUnifiedMem() override;
+
+  void* getDefaultStream() override;
+  void syncDefaultStreamWithHost() override;
+
+  bool isCapableOfGraphCapturing() override;
+  DeviceGraphHandle streamBeginCapture(std::vector<void*>& streamPtrs) override;
+  void streamEndCapture(DeviceGraphHandle handle) override;
+  void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override;
+
+  void* createStream(double priority) override;
+  void destroyGenericStream(void* streamPtr) override;
+  void syncStreamWithHost(void* streamPtr) override;
+  bool isStreamWorkDone(void* streamPtr) override;
+  void syncStreamWithEvent(void* streamPtr, void* eventPtr) override;
+  void streamHostFunction(void* streamPtr, const std::function<void()>& function) override;
+
+  void streamWaitMemory(void* streamPtr, uint32_t* location, uint32_t value) override;
+
+  void* createEvent(bool withTiming) override;
+  void destroyEvent(void* eventPtr) override;
+  void syncEventWithHost(void* eventPtr) override;
+  bool isEventCompleted(void* eventPtr) override;
+  void recordEventOnHost(void* eventPtr) override;
+  void recordEventOnStream(void* eventPtr, void* streamPtr) override;
+  double timespanEvents(void* eventPtrStart, void* eventPtrEnd) override;
+
+  bool isUnifiedMemoryDefault() override;
+
+  void initialize() override;
+  void finalize() override;
+  void profilingMessage(const std::string& message) override;
+  void putProfilingMark(const std::string& name, ProfilingColors color) override;
+  void popLastProfilingMark() override;
+
+  void setupPrinting(int rank) override;
+
+  private:
+  // a stream only needs an identity
+  struct Stream {};
+
+  struct Event {
+    std::chrono::steady_clock::time_point recorded{};
+  };
+
+  void* allocate(size_t size, bool unified);
+  void deallocate(void* ptr);
+
+  static Event* toEvent(void* eventPtr);
+
+  device::StatusT status{false};
+
+  Stream defaultStream{};
+  std::unordered_map<void*, std::unique_ptr<Stream>> genericStreams{};
+
+  // the same bookkeeping as in the other backends
+  size_t allocatedMemBytes{0};
+  size_t allocatedUnifiedMemBytes{0};
+  size_t deallocatedMemBytes{0};
+  std::unordered_map<void*, size_t> memToSizeMap{};
+};
+} // namespace device
+
+#endif // SEISSOLDEVICE_INTERFACES_HOST_HOSTWRAPPEDAPI_H_
diff --git a/interfaces/host/Memory.cpp b/interfaces/host/Memory.cpp
new file mode 100644
index 0000000..e1f12fd
--- /dev/null
+++ b/interfaces/host/Memory.cpp
@@ -0,0 +1,146 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "HostWrappedAPI.h"
+#include "utils/logger.h"
+
+#include <cstdlib>
+#include <cstring>
+#include <mutex>
+#include <sstream>
+#include <string>
+
+using namespace device;
+
+void* ConcreteAPI::allocate(size_t size, bool unified) {
+  isFlagSet<DeviceSelected>(status);
+  if (size == 0) {
+    return nullptr;
+  }
+  const size_t alignment = getGlobMemAlignment();
+  // std::aligned_alloc requires the size to be a multiple of the alignment
+  const size_t paddedSize = ((size + alignment - 1) / alignment) * alignment;
+  void* ptr = std::aligned_alloc(alignment, paddedSize);
+  if (ptr == nullptr) {
+    logError() << "The host backend failed to allocate" << size << "bytes";
+    return nullptr;
+  }
+
+  std::lock_guard guard(apiMutex);
+  allocatedMemBytes += size;
+  if (unified) {
+    allocatedUnifiedMemBytes += size;
+  }
+  memToSizeMap[ptr] = size;
+  return ptr;
+}
+
+void ConcreteAPI::deallocate(void* ptr) {
+  isFlagSet<DeviceSelected>(status);
+  if (ptr == nullptr) {
+    return;
+  }
+  std::lock_guard guard(apiMutex);
+  const auto it = memToSizeMap.find(ptr);
+  if (it == memToSizeMap.end()) {
+    logError() << "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer";
+    return;
+  }
+  deallocatedMemBytes += it->second;
+  memToSizeMap.erase(it);
+  std::free(ptr);
+}
+
+void* ConcreteAPI::allocGlobMem(size_t size, bool /*compress*/) { return allocate(size, false); }
+
+void* ConcreteAPI::allocUnifiedMem(size_t size, bool /*compress*/, Destination /*hint*/) {
+  return allocate(size, true);
+}
+
+void* ConcreteAPI::allocPinnedMem(size_t size, bool /*compress*/, Destination /*hint*/) {
+  return allocate(size, false);
+}
+
+void ConcreteAPI::freeGlobMem(void* devPtr) { deallocate(devPtr); }
+
+void ConcreteAPI::freeUnifiedMem(void* devPtr) { deallocate(devPtr); }
+
+void ConcreteAPI::freePinnedMem(void* devPtr) { deallocate(devPtr); }
+
+std::string ConcreteAPI::getMemLeaksReport() {
+  isFlagSet<DeviceSelected>(status);
+  std::lock_guard guard(apiMutex);
+  std::ostringstream report{};
+  report << "Memory Leaks, bytes: " << (allocatedMemBytes - deallocatedMemBytes) << '\n';
+  return report.str();
+}
+
+size_t ConcreteAPI::getCurrentlyOccupiedMem() {
+  isFlagSet<DeviceSelected>(status);
+  std::lock_guard guard(apiMutex);
+  return allocatedMemBytes;
+}
+
+size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() {
+  isFlagSet<DeviceSelected>(status);
+  std::lock_guard guard(apiMutex);
+  return allocatedUnifiedMemBytes;
+}
+
+void* ConcreteAPI::allocMemAsync(size_t size, void* /*streamPtr*/) { return allocate(size, false); }
+
+void ConcreteAPI::freeMemAsync(void* devPtr, void* /*streamPtr*/) { deallocate(devPtr); }
+
+void ConcreteAPI::pinMemory(void* /*ptr*/, size_t /*size*/) {}
+
+void ConcreteAPI::unpinMemory(void* /*ptr*/) {}
+
+void* ConcreteAPI::devicePointer(void* ptr) { return ptr; }
+
+void ConcreteAPI::copyTo(void* dst, const void* src, size_t count) { std::memcpy(dst, src, count); }
+
+void ConcreteAPI::copyFrom(void* dst, const void* src, size_t count) {
+  std::memcpy(dst, src, count);
+}
+
+void ConcreteAPI::copyBetween(void* dst, const void* src, size_t count) {
+  std::memcpy(dst, src, count);
+}
+
+void ConcreteAPI::copyToAsync(void* dst, const void* src, size_t count, void* /*streamPtr*/) {
+  std::memcpy(dst, src, count);
+}
+
+void ConcreteAPI::copyFromAsync(void* dst, const void* src, size_t count, void* /*streamPtr*/) {
+  std::memcpy(dst, src, count);
+}
+
+void ConcreteAPI::copyBetweenAsync(void* dst, const void* src, size_t count, void* /*streamPtr*/) {
+  std::memcpy(dst, src, count);
+}
+
+namespace {
+void copy2d(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height) {
+  auto* dstBytes = static_cast<char*>(dst);
+  const auto* srcBytes = static_cast<const char*>(src);
+  for (size_t row = 0; row < height; ++row) {
+    std::memcpy(dstBytes + row * dpitch, srcBytes + row * spitch, width);
+  }
+}
+} // namespace
+
+void ConcreteAPI::copy2dArrayTo(
+    void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height) {
+  copy2d(dst, dpitch, src, spitch, width, height);
+}
+
+void ConcreteAPI::copy2dArrayFrom(
+    void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height) {
+  copy2d(dst, dpitch, src, spitch, width, height);
+}
+
+void ConcreteAPI::prefetchUnifiedMemTo(Destination /*type*/,
+                                       const void* /*devPtr*/,
+                                       size_t /*count*/,
+                                       void* /*streamPtr*/) {}
diff --git a/interfaces/host/Streams.cpp b/interfaces/host/Streams.cpp
new file mode 100644
index 0000000..185511c
--- /dev/null
+++ b/interfaces/host/Streams.cpp
@@ -0,0 +1,96 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "DataTypes.h"
+#include "HostWrappedAPI.h"
+#include "utils/logger.h"
+
+#include <chrono>
+#include <cstdint>
+#include <functional>
+#include <memory>
+#include <mutex>
+#include <thread>
+#include <vector>
+
+using namespace device;
+
+ConcreteAPI::Event* ConcreteAPI::toEvent(void* eventPtr) { return static_cast<Event*>(eventPtr); }
+
+void* ConcreteAPI::getDefaultStream() {
+  isFlagSet<InterfaceInitialized>(status);
+  return &defaultStream;
+}
+
+void ConcreteAPI::syncDefaultStreamWithHost() { isFlagSet<InterfaceInitialized>(status); }
+
+void* ConcreteAPI::createStream(double /*priority*/) {
+  isFlagSet<InterfaceInitialized>(status);
+  auto stream = std::make_unique<Stream>();
+  void* streamPtr = stream.get();
+  std::lock_guard guard(apiMutex);
+  genericStreams.emplace(streamPtr, std::move(stream));
+  return streamPtr;
+}
+
+void ConcreteAPI::destroyGenericStream(void* streamPtr) {
+  isFlagSet<InterfaceInitialized>(status);
+  std::lock_guard guard(apiMutex);
+  genericStreams.erase(streamPtr);
+}
+
+void ConcreteAPI::syncStreamWithHost(void* /*streamPtr*/) {
+  isFlagSet<InterfaceInitialized>(status);
+}
+
+bool ConcreteAPI::isStreamWorkDone(void* /*streamPtr*/) {
+  isFlagSet<InterfaceInitialized>(status);
+  return true;
+}
+
+void ConcreteAPI::syncStreamWithEvent(void* /*streamPtr*/, void* /*eventPtr*/) {}
+
+void ConcreteAPI::streamHostFunction(void* /*streamPtr*/, const std::function<void()>& function) {
+  // everything enqueued before has completed already; thus, the function can run right away
+  function();
+}
+
+void ConcreteAPI::streamWaitMemory(void* /*streamPtr*/, uint32_t* location, uint32_t value) {
+  // the stream runs on the calling thread; hence, it is the caller that waits until another
+  // thread has written the value
+  while (__atomic_load_n(location, __ATOMIC_ACQUIRE) < value) {
+    std::this_thread::yield();
+  }
+}
+
+bool ConcreteAPI::isCapableOfGraphCapturing() { return false; }
+
+DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector<void*>& /*streamPtrs*/) {
+  return DeviceGraphHandle();
+}
+
+void ConcreteAPI::streamEndCapture(DeviceGraphHandle /*handle*/) {}
+
+void ConcreteAPI::launchGraph(DeviceGraphHandle /*graphHandle*/, void* /*streamPtr*/) {}
+
+void* ConcreteAPI::createEvent(bool /*withTiming*/) { return new Event(); }
+
+void ConcreteAPI::destroyEvent(void* eventPtr) { delete toEvent(eventPtr); }
+
+void ConcreteAPI::syncEventWithHost(void* /*eventPtr*/) {}
+
+bool ConcreteAPI::isEventCompleted(void* /*eventPtr*/) { return true; }
+
+void ConcreteAPI::recordEventOnHost(void* eventPtr) {
+  toEvent(eventPtr)->recorded = std::chrono::steady_clock::now();
+}
+
+void ConcreteAPI::recordEventOnStream(void* eventPtr, void* /*streamPtr*/) {
+  toEvent(eventPtr)->recorded = std::chrono::steady_clock::now();
+}
+
+double ConcreteAPI::timespanEvents(void* eventPtrStart, void* eventPtrEnd) {
+  const auto span = toEvent(eventPtrEnd)->recorded - toEvent(eventPtrStart)->recorded;
+  return std::chrono::duration<double>(span).count();
+}
diff --git a/tests/BaseTestSuite.h b/tests/BaseTestSuite.h
index 4d2c2ae..5ad8733 100644
--- a/tests/BaseTestSuite.h
+++ b/tests/BaseTestSuite.h
@@ -23,7 +23,7 @@ class BaseTestSuite : public ::testing::Test {
   BaseTestSuite() { randomEngine.seed(randomDevice()); }
 
   void SetUp() {
-    device = &DeviceInstance::getInstance();
+    device = &DeviceInstance::instance();
     setUp = true;
   }
 
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index ece56b8..7c7b99a 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -16,7 +16,7 @@ set_property(CACHE SM PROPERTY STRINGS ${SM_OPTIONS})
 set(DEVICE_ARCH ${SM})
 
 set(DEVICE_BACKEND "cuda" CACHE STRING "type of an interface")
-set_property(CACHE DEVICE_BACKEND PROPERTY STRINGS "cuda" "hip" "oneapi" "hipsycl")
+set_property(CACHE DEVICE_BACKEND PROPERTY STRINGS "none" "cuda" "hip" "oneapi" "hipsycl" "acpp")
 
 add_subdirectory(.. root)
 
@@ -31,5 +31,12 @@ if (DEVICE_BACKEND STREQUAL "hipsycl")
   target_link_libraries(tests PRIVATE OpenMP::OpenMP_CXX)
 endif()
 
+if (DEVICE_BACKEND STREQUAL "none")
+  # the host backend guarantees more than the common API; these tests check that part
+  target_sources(tests PRIVATE host.cpp)
+  find_package(Threads REQUIRED)
+  target_link_libraries(tests PRIVATE Threads::Threads)
+endif()
+
 enable_testing()
 add_test(NAME all_tests COMMAND tests)
diff --git a/tests/array_manip.cpp b/tests/array_manip.cpp
index 71df856..1bd1741 100644
--- a/tests/array_manip.cpp
+++ b/tests/array_manip.cpp
@@ -17,60 +17,62 @@ class ArrayManip : public BaseTestSuite {
 
 TEST_F(ArrayManip, fill) {
   const int N = 100;
-  int* arr = (int*)device->api->allocGlobMem(N * sizeof(int));
+  int* arr = (int*)device->api().allocGlobMem(N * sizeof(int));
   int scalar = 502;
 
-  device->algorithms.fillArray(arr, scalar, N, device->api->getDefaultStream());
+  device->algorithms().fillArray(arr, scalar, N, device->api().getDefaultStream());
 
   std::vector<int> hostVector(N, 0);
-  device->api->copyFromAsync(&hostVector[0], arr, N * sizeof(int), device->api->getDefaultStream());
+  device->api().copyFromAsync(
+      &hostVector[0], arr, N * sizeof(int), device->api().getDefaultStream());
 
-  device->api->syncDefaultStreamWithHost();
+  device->api().syncDefaultStreamWithHost();
 
   for (auto& i : hostVector) {
     EXPECT_EQ(scalar, i);
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 
 TEST_F(ArrayManip, touchClean32) {
 
   const int N = 100;
-  float* arr = (float*)device->api->allocGlobMem(N * sizeof(float));
-  device->algorithms.touchMemory(arr, N, true, device->api->getDefaultStream());
+  float* arr = (float*)device->api().allocGlobMem(N * sizeof(float));
+  device->algorithms().touchMemory(arr, N, true, device->api().getDefaultStream());
   std::vector<float> hostVector(N, 1);
 
-  device->api->copyFromAsync(
-      &hostVector[0], arr, N * sizeof(float), device->api->getDefaultStream());
+  device->api().copyFromAsync(
+      &hostVector[0], arr, N * sizeof(float), device->api().getDefaultStream());
 
-  device->api->syncDefaultStreamWithHost();
+  device->api().syncDefaultStreamWithHost();
 
   for (auto& i : hostVector) {
     EXPECT_EQ(0, i);
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 
 TEST_F(ArrayManip, touchNoClean32) {
 
   const int N = 100;
-  float* arr = (float*)device->api->allocGlobMem(N * sizeof(float));
+  float* arr = (float*)device->api().allocGlobMem(N * sizeof(float));
   std::vector<float> hostVector(N, 0);
 
-  device->api->copyToAsync(arr, &hostVector[0], N * sizeof(float), device->api->getDefaultStream());
-  device->algorithms.touchMemory(arr, N, false, device->api->getDefaultStream());
-  device->api->copyFromAsync(
-      &hostVector[0], arr, N * sizeof(float), device->api->getDefaultStream());
+  device->api().copyToAsync(
+      arr, &hostVector[0], N * sizeof(float), device->api().getDefaultStream());
+  device->algorithms().touchMemory(arr, N, false, device->api().getDefaultStream());
+  device->api().copyFromAsync(
+      &hostVector[0], arr, N * sizeof(float), device->api().getDefaultStream());
 
-  device->api->syncDefaultStreamWithHost();
+  device->api().syncDefaultStreamWithHost();
 
   for (auto& i : hostVector) {
     EXPECT_EQ(0, i);
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 
 // avoid double, due to some archs (also in the CI) don't seem to thoroughly support it
@@ -78,56 +80,58 @@ TEST_F(ArrayManip, touchNoClean32) {
 TEST_F(ArrayManip, touchClean64) {
 
   const int N = 100;
-  long* arr = (long*)device->api->allocGlobMem(N * sizeof(long));
-  device->algorithms.touchMemory(arr, N, true, device->api->getDefaultStream());
+  long* arr = (long*)device->api().allocGlobMem(N * sizeof(long));
+  device->algorithms().touchMemory(arr, N, true, device->api().getDefaultStream());
   std::vector<long> hostVector(N, 1);
 
-  device->api->copyFromAsync(
-      &hostVector[0], arr, N * sizeof(long), device->api->getDefaultStream());
+  device->api().copyFromAsync(
+      &hostVector[0], arr, N * sizeof(long), device->api().getDefaultStream());
 
-  device->api->syncDefaultStreamWithHost();
+  device->api().syncDefaultStreamWithHost();
 
   for (auto& i : hostVector) {
     EXPECT_EQ(0, i);
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 
 TEST_F(ArrayManip, touchNoClean64) {
 
   const int N = 100;
-  long* arr = (long*)device->api->allocGlobMem(N * sizeof(long));
+  long* arr = (long*)device->api().allocGlobMem(N * sizeof(long));
   std::vector<long> hostVector(N, 0);
 
-  device->api->copyToAsync(arr, &hostVector[0], N * sizeof(long), device->api->getDefaultStream());
-  device->algorithms.touchMemory(arr, N, false, device->api->getDefaultStream());
-  device->api->copyFromAsync(
-      &hostVector[0], arr, N * sizeof(long), device->api->getDefaultStream());
+  device->api().copyToAsync(
+      arr, &hostVector[0], N * sizeof(long), device->api().getDefaultStream());
+  device->algorithms().touchMemory(arr, N, false, device->api().getDefaultStream());
+  device->api().copyFromAsync(
+      &hostVector[0], arr, N * sizeof(long), device->api().getDefaultStream());
 
-  device->api->syncDefaultStreamWithHost();
+  device->api().syncDefaultStreamWithHost();
 
   for (auto& i : hostVector) {
     EXPECT_EQ(0, i);
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 
 TEST_F(ArrayManip, scale) {
   const int N = 100;
   std::vector<int> hostVector(N, 1);
-  int* arr = (int*)device->api->allocGlobMem(N * sizeof(int));
+  int* arr = (int*)device->api().allocGlobMem(N * sizeof(int));
 
-  device->api->copyToAsync(arr, &hostVector[0], N * sizeof(int), device->api->getDefaultStream());
-  device->algorithms.scaleArray(arr, 5, N, device->api->getDefaultStream());
-  device->api->copyFromAsync(&hostVector[0], arr, N * sizeof(int), device->api->getDefaultStream());
+  device->api().copyToAsync(arr, &hostVector[0], N * sizeof(int), device->api().getDefaultStream());
+  device->algorithms().scaleArray(arr, 5, N, device->api().getDefaultStream());
+  device->api().copyFromAsync(
+      &hostVector[0], arr, N * sizeof(int), device->api().getDefaultStream());
 
-  device->api->syncDefaultStreamWithHost();
+  device->api().syncDefaultStreamWithHost();
 
   for (auto& i : hostVector) {
     EXPECT_EQ(5, i);
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
diff --git a/tests/batch_manip.cpp b/tests/batch_manip.cpp
index 54ad1b1..24ff18e 100644
--- a/tests/batch_manip.cpp
+++ b/tests/batch_manip.cpp
@@ -18,8 +18,8 @@ class BatchManip : public BaseTestSuite {
   public:
   template <typename T, typename F>
   void testWrapper(size_t N, size_t M, bool sparse, F&& inner) {
-    T* data = (T*)device->api->allocGlobMem(N * M * sizeof(T));
-    T** batch = (T**)device->api->allocUnifiedMem(N * sizeof(T*));
+    T* data = (T*)device->api().allocGlobMem(N * M * sizeof(T));
+    T** batch = (T**)device->api().allocUnifiedMem(N * sizeof(T*));
 
     for (size_t i = 0; i < N; ++i) {
       if (!(sparse && i % 2 == 0)) {
@@ -31,8 +31,8 @@ class BatchManip : public BaseTestSuite {
 
     std::forward<F>(inner)(batch, data);
 
-    device->api->freeGlobMem(data);
-    device->api->freeUnifiedMem(batch);
+    device->api().freeGlobMem(data);
+    device->api().freeUnifiedMem(batch);
   }
 };
 
@@ -42,13 +42,13 @@ TEST_F(BatchManip, fill32) {
   testWrapper<float>(N, M, false, [&](float** batch, float* data) {
     float scalar = 502;
 
-    device->algorithms.setToValue(batch, scalar, M, N, device->api->getDefaultStream());
+    device->algorithms().setToValue(batch, scalar, M, N, device->api().getDefaultStream());
 
     std::vector<float> hostVector(N * M, 0);
-    device->api->copyFromAsync(
-        &hostVector[0], data, N * M * sizeof(float), device->api->getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, N * M * sizeof(float), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(scalar, i);
@@ -60,13 +60,13 @@ TEST_F(BatchManip, touchClean32) {
   const int N = 100;
   const int M = 120;
   testWrapper<float>(N, M, false, [&](float** batch, float* data) {
-    device->algorithms.touchBatchedMemory(batch, M, N, true, device->api->getDefaultStream());
+    device->algorithms().touchBatchedMemory(batch, M, N, true, device->api().getDefaultStream());
     std::vector<float> hostVector(N * M, 1);
 
-    device->api->copyFromAsync(
-        &hostVector[0], data, M * N * sizeof(float), device->api->getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, M * N * sizeof(float), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(0, i);
@@ -76,13 +76,13 @@ TEST_F(BatchManip, touchClean32) {
   testWrapper<float>(N, M, true, [&](float** batch, float* data) {
     std::vector<float> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data, &hostVector[0], N * M * sizeof(float), device->api->getDefaultStream());
-    device->algorithms.touchBatchedMemory(batch, M, N, true, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data, M * N * sizeof(float), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data, &hostVector[0], N * M * sizeof(float), device->api().getDefaultStream());
+    device->algorithms().touchBatchedMemory(batch, M, N, true, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, M * N * sizeof(float), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (size_t i = 0; i < N; ++i) {
       for (size_t j = 0; j < M; ++j) {
@@ -102,13 +102,13 @@ TEST_F(BatchManip, touchNoClean32) {
   testWrapper<float>(N, M, false, [&](float** batch, float* data) {
     std::vector<float> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data, &hostVector[0], N * M * sizeof(float), device->api->getDefaultStream());
-    device->algorithms.touchBatchedMemory(batch, M, N, false, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data, N * M * sizeof(float), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data, &hostVector[0], N * M * sizeof(float), device->api().getDefaultStream());
+    device->algorithms().touchBatchedMemory(batch, M, N, false, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, N * M * sizeof(float), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(1, i);
@@ -120,47 +120,48 @@ TEST_F(BatchManip, scatterToUniform32) {
   const int N = 10000;
   const int M = 12000;
 
-  float* data2 = (float*)device->api->allocGlobMem(N * M * sizeof(float));
+  float* data2 = (float*)device->api().allocGlobMem(N * M * sizeof(float));
   testWrapper<float>(N, M, false, [&](float** batch, float* data) {
     std::vector<float> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data, &hostVector[0], N * M * sizeof(float), device->api->getDefaultStream());
-    device->algorithms.copyScatterToUniform(
-        const_cast<const float**>(batch), data2, M, M, N, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data2, N * M * sizeof(float), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data, &hostVector[0], N * M * sizeof(float), device->api().getDefaultStream());
+    device->algorithms().copyScatterToUniform(
+        const_cast<const float**>(batch), data2, M, M, N, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data2, N * M * sizeof(float), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(1, i);
     }
   });
-  device->api->freeGlobMem(data2);
+  device->api().freeGlobMem(data2);
 }
 
 TEST_F(BatchManip, uniformToScatter32) {
   const int N = 10000;
   const int M = 12000;
 
-  float* data2 = (float*)device->api->allocGlobMem(N * M * sizeof(float));
+  float* data2 = (float*)device->api().allocGlobMem(N * M * sizeof(float));
   testWrapper<float>(N, M, false, [&](float** batch, float* data) {
     std::vector<float> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data2, &hostVector[0], N * M * sizeof(float), device->api->getDefaultStream());
-    device->algorithms.copyUniformToScatter(data2, batch, M, M, N, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data, N * M * sizeof(float), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data2, &hostVector[0], N * M * sizeof(float), device->api().getDefaultStream());
+    device->algorithms().copyUniformToScatter(
+        data2, batch, M, M, N, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, N * M * sizeof(float), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(1, i);
     }
   });
-  device->api->freeGlobMem(data2);
+  device->api().freeGlobMem(data2);
 }
 
 TEST_F(BatchManip, fill64) {
@@ -169,13 +170,13 @@ TEST_F(BatchManip, fill64) {
   testWrapper<long>(N, M, false, [&](long** batch, long* data) {
     long scalar = 502;
 
-    device->algorithms.setToValue(batch, scalar, M, N, device->api->getDefaultStream());
+    device->algorithms().setToValue(batch, scalar, M, N, device->api().getDefaultStream());
 
     std::vector<long> hostVector(N * M, 0);
-    device->api->copyFromAsync(
-        &hostVector[0], data, N * M * sizeof(long), device->api->getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, N * M * sizeof(long), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(scalar, i);
@@ -187,13 +188,13 @@ TEST_F(BatchManip, touchClean64) {
   const int N = 100;
   const int M = 120;
   testWrapper<long>(N, M, false, [&](long** batch, long* data) {
-    device->algorithms.touchBatchedMemory(batch, M, N, true, device->api->getDefaultStream());
+    device->algorithms().touchBatchedMemory(batch, M, N, true, device->api().getDefaultStream());
     std::vector<long> hostVector(N * M, 1);
 
-    device->api->copyFromAsync(
-        &hostVector[0], data, M * N * sizeof(long), device->api->getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, M * N * sizeof(long), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(0, i);
@@ -203,13 +204,13 @@ TEST_F(BatchManip, touchClean64) {
   testWrapper<long>(N, M, true, [&](long** batch, long* data) {
     std::vector<long> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data, &hostVector[0], N * M * sizeof(long), device->api->getDefaultStream());
-    device->algorithms.touchBatchedMemory(batch, M, N, true, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data, M * N * sizeof(long), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data, &hostVector[0], N * M * sizeof(long), device->api().getDefaultStream());
+    device->algorithms().touchBatchedMemory(batch, M, N, true, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, M * N * sizeof(long), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (size_t i = 0; i < N; ++i) {
       for (size_t j = 0; j < M; ++j) {
@@ -229,13 +230,13 @@ TEST_F(BatchManip, touchNoClean64) {
   testWrapper<long>(N, M, false, [&](long** batch, long* data) {
     std::vector<long> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data, &hostVector[0], N * M * sizeof(long), device->api->getDefaultStream());
-    device->algorithms.touchBatchedMemory(batch, M, N, false, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data, N * M * sizeof(long), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data, &hostVector[0], N * M * sizeof(long), device->api().getDefaultStream());
+    device->algorithms().touchBatchedMemory(batch, M, N, false, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, N * M * sizeof(long), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(1, i);
@@ -247,45 +248,46 @@ TEST_F(BatchManip, scatterToUniform64) {
   const int N = 100;
   const int M = 120;
 
-  long* data2 = (long*)device->api->allocGlobMem(N * M * sizeof(long));
+  long* data2 = (long*)device->api().allocGlobMem(N * M * sizeof(long));
   testWrapper<long>(N, M, false, [&](long** batch, long* data) {
     std::vector<long> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data, &hostVector[0], N * M * sizeof(long), device->api->getDefaultStream());
-    device->algorithms.copyScatterToUniform(
-        const_cast<const long**>(batch), data2, M, M, N, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data2, N * M * sizeof(long), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data, &hostVector[0], N * M * sizeof(long), device->api().getDefaultStream());
+    device->algorithms().copyScatterToUniform(
+        const_cast<const long**>(batch), data2, M, M, N, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data2, N * M * sizeof(long), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(1, i);
     }
   });
-  device->api->freeGlobMem(data2);
+  device->api().freeGlobMem(data2);
 }
 
 TEST_F(BatchManip, uniformToScatter64) {
   const int N = 100;
   const int M = 120;
 
-  long* data2 = (long*)device->api->allocGlobMem(N * M * sizeof(long));
+  long* data2 = (long*)device->api().allocGlobMem(N * M * sizeof(long));
   testWrapper<long>(N, M, false, [&](long** batch, long* data) {
     std::vector<long> hostVector(N * M, 1);
 
-    device->api->copyToAsync(
-        data2, &hostVector[0], N * M * sizeof(long), device->api->getDefaultStream());
-    device->algorithms.copyUniformToScatter(data2, batch, M, M, N, device->api->getDefaultStream());
-    device->api->copyFromAsync(
-        &hostVector[0], data, N * M * sizeof(long), device->api->getDefaultStream());
+    device->api().copyToAsync(
+        data2, &hostVector[0], N * M * sizeof(long), device->api().getDefaultStream());
+    device->algorithms().copyUniformToScatter(
+        data2, batch, M, M, N, device->api().getDefaultStream());
+    device->api().copyFromAsync(
+        &hostVector[0], data, N * M * sizeof(long), device->api().getDefaultStream());
 
-    device->api->syncDefaultStreamWithHost();
+    device->api().syncDefaultStreamWithHost();
 
     for (auto& i : hostVector) {
       EXPECT_EQ(1, i);
     }
   });
-  device->api->freeGlobMem(data2);
+  device->api().freeGlobMem(data2);
 }
diff --git a/tests/host.cpp b/tests/host.cpp
new file mode 100644
index 0000000..32c88db
--- /dev/null
+++ b/tests/host.cpp
@@ -0,0 +1,193 @@
+// SPDX-FileCopyrightText: 2026 SeisSol Group
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "BaseTestSuite.h"
+#include "device.h"
+
+#include "gtest/gtest.h"
+#include <chrono>
+#include <cstdint>
+#include <cstring>
+#include <string>
+#include <thread>
+#include <vector>
+
+using namespace device;
+using namespace ::testing;
+
+// the behavior that the host backend (DEVICE_BACKEND=none) guarantees on top of the common API
+class Host : public BaseTestSuite {
+  using BaseTestSuite::BaseTestSuite;
+};
+
+namespace {
+size_t leakedBytes(AbstractAPI& api) {
+  const std::string report = api.getMemLeaksReport();
+  const std::string prefix = "bytes: ";
+  return std::stoul(report.substr(report.find(prefix) + prefix.size()));
+}
+} // namespace
+
+TEST_F(Host, properties) {
+  auto& api = device->api();
+  EXPECT_EQ(api.getApiName(), "Host");
+  EXPECT_EQ(api.getNumDevices(), 1);
+  EXPECT_EQ(api.getDeviceId(), 0);
+  EXPECT_FALSE(api.isCapableOfGraphCapturing());
+  EXPECT_TRUE(api.isUnifiedMemoryDefault());
+  EXPECT_GT(api.getMaxAvailableMem(), 0);
+}
+
+TEST_F(Host, allocationsAreAligned) {
+  auto& api = device->api();
+  const auto alignment = api.getGlobMemAlignment();
+  for (const size_t size : {1, 7, 128, 1000}) {
+    void* glob = api.allocGlobMem(size);
+    void* unified = api.allocUnifiedMem(size);
+    void* pinned = api.allocPinnedMem(size);
+    for (void* ptr : {glob, unified, pinned}) {
+      ASSERT_NE(ptr, nullptr);
+      EXPECT_EQ(reinterpret_cast<uintptr_t>(ptr) % alignment, 0);
+    }
+    api.freeGlobMem(glob);
+    api.freeUnifiedMem(unified);
+    api.freePinnedMem(pinned);
+  }
+  EXPECT_EQ(api.allocGlobMem(0), nullptr);
+  api.freeGlobMem(nullptr);
+}
+
+TEST_F(Host, bookkeeping) {
+  auto& api = device->api();
+  const auto occupied = api.getCurrentlyOccupiedMem();
+  const auto occupiedUnified = api.getCurrentlyOccupiedUnifiedMem();
+  const auto leaked = leakedBytes(api);
+
+  void* glob = api.allocGlobMem(100);
+  void* unified = api.allocUnifiedMem(50);
+  EXPECT_EQ(api.getCurrentlyOccupiedMem(), occupied + 150);
+  EXPECT_EQ(api.getCurrentlyOccupiedUnifiedMem(), occupiedUnified + 50);
+  EXPECT_EQ(leakedBytes(api), leaked + 150);
+
+  api.freeGlobMem(glob);
+  api.freeUnifiedMem(unified);
+  EXPECT_EQ(leakedBytes(api), leaked);
+}
+
+TEST_F(Host, hostFunctionsRunInOrder) {
+  auto& api = device->api();
+  void* stream = api.createStream();
+  std::vector<int> order;
+  api.streamHostFunction(stream, [&] { order.push_back(1); });
+  api.streamHostFunction(stream, [&] { order.push_back(2); });
+  api.syncStreamWithHost(stream);
+  EXPECT_EQ(order, (std::vector<int>{1, 2}));
+  EXPECT_TRUE(api.isStreamWorkDone(stream));
+  api.destroyGenericStream(stream);
+}
+
+TEST_F(Host, eventsMeasureTime) {
+  auto& api = device->api();
+  void* stream = api.getDefaultStream();
+  void* start = api.createEvent(true);
+  void* end = api.createEvent(true);
+  api.recordEventOnStream(start, stream);
+  std::this_thread::sleep_for(std::chrono::milliseconds(20));
+  api.recordEventOnStream(end, stream);
+  api.syncEventWithHost(end);
+  EXPECT_TRUE(api.isEventCompleted(end));
+  EXPECT_GE(api.timespanEvents(start, end), 0.019);
+  api.destroyEvent(start);
+  api.destroyEvent(end);
+}
+
+TEST_F(Host, waitMemory) {
+  auto& api = device->api();
+  auto* flag = static_cast<uint32_t*>(
+      api.allocPinnedMem(sizeof(uint32_t), false, Destination::CurrentDevice));
+
+  // already reached: returns right away
+  __atomic_store_n(flag, 5U, __ATOMIC_RELEASE);
+  api.streamWaitMemory(api.getDefaultStream(), flag, 3);
+
+  // reached later, by another thread
+  __atomic_store_n(flag, 0U, __ATOMIC_RELEASE);
+  std::thread writer([flag] {
+    std::this_thread::sleep_for(std::chrono::milliseconds(10));
+    __atomic_store_n(flag, 2U, __ATOMIC_RELEASE);
+  });
+  api.streamWaitMemory(api.getDefaultStream(), flag, 2);
+  EXPECT_EQ(__atomic_load_n(flag, __ATOMIC_ACQUIRE), 2U);
+  writer.join();
+
+  api.freePinnedMem(flag);
+}
+
+TEST_F(Host, graphCapturingIsUnavailable) {
+  auto& api = device->api();
+  std::vector<void*> streams{api.getDefaultStream()};
+  const auto handle = api.streamBeginCapture(streams);
+  EXPECT_FALSE(handle.isInitialized());
+  api.streamEndCapture(handle);
+}
+
+TEST_F(Host, reductions) {
+  auto& api = device->api();
+  const std::vector<float> values{-3.0F, -1.5F, -7.25F};
+  const auto bytes = values.size() * sizeof(float);
+  auto* buffer = static_cast<float*>(api.allocGlobMem(bytes));
+  api.copyTo(buffer, values.data(), bytes);
+  auto* result = static_cast<float*>(api.allocPinnedMem(sizeof(float)));
+  void* stream = api.getDefaultStream();
+
+  // the neutral element of the maximum is the lowest value, not the smallest positive one
+  device->algorithms().reduceVector(
+      result, buffer, true, values.size(), ReductionType::Max, stream);
+  EXPECT_EQ(*result, -1.5F);
+
+  device->algorithms().reduceVector(
+      result, buffer, true, values.size(), ReductionType::Min, stream);
+  EXPECT_EQ(*result, -7.25F);
+
+  // without overriding, the reduction continues from the value in the result
+  *result = 10.0F;
+  device->algorithms().reduceVector(
+      result, buffer, false, values.size(), ReductionType::Add, stream);
+  EXPECT_EQ(*result, -1.75F);
+
+  api.freePinnedMem(result);
+  api.freeGlobMem(buffer);
+}
+
+TEST_F(Host, pitchedCopies) {
+  auto& api = device->api();
+  constexpr size_t Width = 3;
+  constexpr size_t Height = 4;
+  constexpr size_t SrcPitch = 5;
+  constexpr size_t DevPitch = 7;
+  constexpr char Padding = '#';
+
+  std::vector<char> src(SrcPitch * Height, Padding);
+  for (size_t row = 0; row < Height; ++row) {
+    for (size_t column = 0; column < Width; ++column) {
+      src[row * SrcPitch + column] = static_cast<char>('a' + row * Width + column);
+    }
+  }
+
+  auto* dev = static_cast<char*>(api.allocGlobMem(DevPitch * Height));
+  std::memset(dev, Padding, DevPitch * Height);
+  api.copy2dArrayTo(dev, DevPitch, src.data(), SrcPitch, Width, Height);
+
+  std::vector<char> dst(SrcPitch * Height, Padding);
+  api.copy2dArrayFrom(dst.data(), SrcPitch, dev, DevPitch, Width, Height);
+  EXPECT_EQ(dst, src);
+
+  for (size_t row = 0; row < Height; ++row) {
+    for (size_t column = Width; column < DevPitch; ++column) {
+      EXPECT_EQ(dev[row * DevPitch + column], Padding);
+    }
+  }
+
+  api.freeGlobMem(dev);
+}
diff --git a/tests/main.cpp b/tests/main.cpp
index cfba087..ac224f6 100644
--- a/tests/main.cpp
+++ b/tests/main.cpp
@@ -10,14 +10,14 @@ using namespace device;
 int main(int argc, char** argv) {
   ::testing::InitGoogleTest(&argc, argv);
 
-  DeviceInstance& device = DeviceInstance::getInstance();
-  device.api->setDevice(0);
-  device.api->initialize();
+  DeviceInstance& device = DeviceInstance::instance();
+  device.api().setDevice(0);
+  device.api().initialize();
 
   const auto result = RUN_ALL_TESTS();
 
-  device.api->syncDevice();
-  device.api->finalize();
+  device.api().syncDevice();
+  device.api().finalize();
 
   return result;
 }
diff --git a/tests/memory.cpp b/tests/memory.cpp
index f4f6a4f..7d483e2 100644
--- a/tests/memory.cpp
+++ b/tests/memory.cpp
@@ -29,16 +29,16 @@ TEST_F(Memories, copy2DMemory) {
     }
   }
 
-  int* arr = (int*)device->api->allocGlobMem(M * N * sizeof(int));
+  int* arr = (int*)device->api().allocGlobMem(M * N * sizeof(int));
 
   int spitch = N * sizeof(int);
   int dpitch = N * sizeof(int);
   int width = N * sizeof(int);
   int height = M;
 
-  device->api->copy2dArrayTo(arr, dpitch, &hostVector[0], spitch, width, height);
+  device->api().copy2dArrayTo(arr, dpitch, &hostVector[0], spitch, width, height);
   int hostVector2[M][N];
-  device->api->copy2dArrayFrom(&hostVector2[0], dpitch, arr, spitch, width, height);
+  device->api().copy2dArrayFrom(&hostVector2[0], dpitch, arr, spitch, width, height);
 
   for (size_t i = 0; i < M; i++) {
     for (size_t j = 0; j < N; j++) {
@@ -46,7 +46,7 @@ TEST_F(Memories, copy2DMemory) {
     }
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 
 /*
@@ -65,16 +65,16 @@ TEST_F(Memories, copy2DMemoryWithSrcPitch) {
     }
   }
 
-  int *arr = (int *)device->api->allocGlobMem(M * N * sizeof(int));
+  int *arr = (int *)device->api().allocGlobMem(M * N * sizeof(int));
 
   int spitch = (N + SPI) * sizeof(int);
   int dpitch = N * sizeof(int);
   int width = N * sizeof(int);
   int height = M;
 
-  device->api->copy2dArrayTo(arr, dpitch, &hostVector[0], spitch, width, height);
+  device->api().copy2dArrayTo(arr, dpitch, &hostVector[0], spitch, width, height);
   int hostVector2[M][N];
-  device->api->copy2dArrayFrom(&hostVector2[0], dpitch, arr, dpitch, width, height);
+  device->api().copy2dArrayFrom(&hostVector2[0], dpitch, arr, dpitch, width, height);
 
   for (size_t i = 0; i < M; i++) {
     for (size_t j = 0; j < N; j++) {
@@ -82,7 +82,7 @@ TEST_F(Memories, copy2DMemoryWithSrcPitch) {
     }
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 */
 
@@ -104,16 +104,16 @@ TEST_F(Memories, copy2DMemoryWithDstPitch) {
     }
   }
 
-  int *arr = (int *)device->api->allocGlobMem(M * (N + DPI) * sizeof(int));
+  int *arr = (int *)device->api().allocGlobMem(M * (N + DPI) * sizeof(int));
 
   int spitch = (N + SPI) * sizeof(int);
   int dpitch = (N + DPI) * sizeof(int);
   int width = N * sizeof(int);
   int height = M;
 
-  device->api->copy2dArrayTo(arr, dpitch, &hostVector[0], spitch, width, height);
+  device->api().copy2dArrayTo(arr, dpitch, &hostVector[0], spitch, width, height);
   int hostVector2[M][N + SPI];
-  device->api->copy2dArrayFrom(&hostVector2[0], spitch, arr, dpitch, width, height);
+  device->api().copy2dArrayFrom(&hostVector2[0], spitch, arr, dpitch, width, height);
 
   for (size_t i = 0; i < M; i++) {
     for (size_t j = 0; j < N + SPI; j++) {
@@ -121,6 +121,6 @@ TEST_F(Memories, copy2DMemoryWithDstPitch) {
     }
   }
 
-  device->api->freeGlobMem(arr);
+  device->api().freeGlobMem(arr);
 }
 */
diff --git a/tests/reductions.cpp b/tests/reductions.cpp
index 0aaac48..edeb2d9 100644
--- a/tests/reductions.cpp
+++ b/tests/reductions.cpp
@@ -28,47 +28,51 @@ TEST_F(Reductions, Add) {
     element = distribution(randomEngine);
   }
 
-  auto* devVector = reinterpret_cast<unsigned*>(device->api->allocGlobMem(sizeof(unsigned) * size));
-  device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size);
+  auto* devVector =
+      reinterpret_cast<unsigned*>(device->api().allocGlobMem(sizeof(unsigned) * size));
+  device->api().copyTo(devVector, vector.data(), sizeof(unsigned) * size);
 
   auto expectedResult = std::accumulate(vector.begin(), vector.end(), 0, std::plus<unsigned>());
 
-  unsigned* testResult = reinterpret_cast<unsigned*>(device->api->allocPinnedMem(sizeof(unsigned)));
+  unsigned* testResult =
+      reinterpret_cast<unsigned*>(device->api().allocPinnedMem(sizeof(unsigned)));
 
-  device->algorithms.reduceVector(
-      testResult, devVector, true, size, ReductionType::Add, device->api->getDefaultStream());
-  device->api->syncDefaultStreamWithHost();
+  device->algorithms().reduceVector(
+      testResult, devVector, true, size, ReductionType::Add, device->api().getDefaultStream());
+  device->api().syncDefaultStreamWithHost();
   EXPECT_EQ(expectedResult, *testResult);
-  device->api->freePinnedMem(testResult);
-  device->api->freeGlobMem(devVector);
+  device->api().freePinnedMem(testResult);
+  device->api().freeGlobMem(devVector);
 }
 
 TEST_F(Reductions, Max) {
   constexpr size_t size = 20010000;
   std::vector<unsigned> vector(size, 0);
 
-  auto* devVector = reinterpret_cast<unsigned*>(device->api->allocGlobMem(sizeof(unsigned) * size));
-  device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size);
+  auto* devVector =
+      reinterpret_cast<unsigned*>(device->api().allocGlobMem(sizeof(unsigned) * size));
+  device->api().copyTo(devVector, vector.data(), sizeof(unsigned) * size);
 
   std::uniform_int_distribution<> distribution(10, 100);
   for (auto& element : vector) {
     element = distribution(randomEngine);
   }
 
-  device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size);
+  device->api().copyTo(devVector, vector.data(), sizeof(unsigned) * size);
 
   auto max = [](unsigned a, unsigned b) -> unsigned { return a > b ? a : b; };
   auto initValue = std::numeric_limits<unsigned>::min();
   auto expectedResult = std::accumulate(vector.begin(), vector.end(), initValue, max);
 
-  unsigned* testResult = reinterpret_cast<unsigned*>(device->api->allocPinnedMem(sizeof(unsigned)));
+  unsigned* testResult =
+      reinterpret_cast<unsigned*>(device->api().allocPinnedMem(sizeof(unsigned)));
 
-  device->algorithms.reduceVector(
-      testResult, devVector, true, size, ReductionType::Max, device->api->getDefaultStream());
-  device->api->syncDefaultStreamWithHost();
+  device->algorithms().reduceVector(
+      testResult, devVector, true, size, ReductionType::Max, device->api().getDefaultStream());
+  device->api().syncDefaultStreamWithHost();
   EXPECT_EQ(expectedResult, *testResult);
-  device->api->freePinnedMem(testResult);
-  device->api->freeGlobMem(devVector);
+  device->api().freePinnedMem(testResult);
+  device->api().freeGlobMem(devVector);
 }
 
 TEST_F(Reductions, Min) {
@@ -80,19 +84,21 @@ TEST_F(Reductions, Min) {
     element = distribution(randomEngine);
   }
 
-  auto* devVector = reinterpret_cast<unsigned*>(device->api->allocGlobMem(sizeof(unsigned) * size));
-  device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size);
+  auto* devVector =
+      reinterpret_cast<unsigned*>(device->api().allocGlobMem(sizeof(unsigned) * size));
+  device->api().copyTo(devVector, vector.data(), sizeof(unsigned) * size);
 
   auto min = [](unsigned a, unsigned b) -> unsigned { return a > b ? b : a; };
   auto initValue = std::numeric_limits<unsigned>::max();
   auto expectedResult = std::accumulate(vector.begin(), vector.end(), initValue, min);
 
-  unsigned* testResult = reinterpret_cast<unsigned*>(device->api->allocPinnedMem(sizeof(unsigned)));
+  unsigned* testResult =
+      reinterpret_cast<unsigned*>(device->api().allocPinnedMem(sizeof(unsigned)));
 
-  device->algorithms.reduceVector(
-      testResult, devVector, true, size, ReductionType::Min, device->api->getDefaultStream());
-  device->api->syncDefaultStreamWithHost();
+  device->algorithms().reduceVector(
+      testResult, devVector, true, size, ReductionType::Min, device->api().getDefaultStream());
+  device->api().syncDefaultStreamWithHost();
   EXPECT_EQ(expectedResult, *testResult);
-  device->api->freePinnedMem(testResult);
-  device->api->freeGlobMem(devVector);
+  device->api().freePinnedMem(testResult);
+  device->api().freeGlobMem(devVector);
 }
diff --git a/tests/streams.cpp b/tests/streams.cpp
index f034445..39120b6 100644
--- a/tests/streams.cpp
+++ b/tests/streams.cpp
@@ -20,58 +20,58 @@ TEST_F(Streams, hostFunctionRunsAfterPrecedingWork) {
   // large enough that the copies are still running when the host function is enqueued
   constexpr std::size_t Count = std::size_t{1} << 24;
 
-  auto* source = static_cast<int*>(device->api->allocPinnedMem(Count * sizeof(int)));
-  auto* target = static_cast<int*>(device->api->allocPinnedMem(Count * sizeof(int)));
-  auto* buffer = static_cast<int*>(device->api->allocGlobMem(Count * sizeof(int)));
+  auto* source = static_cast<int*>(device->api().allocPinnedMem(Count * sizeof(int)));
+  auto* target = static_cast<int*>(device->api().allocPinnedMem(Count * sizeof(int)));
+  auto* buffer = static_cast<int*>(device->api().allocGlobMem(Count * sizeof(int)));
   std::fill(source, source + Count, 1904);
   std::fill(target, target + Count, 0);
 
-  void* stream = device->api->createStream();
-  device->api->copyToAsync(buffer, source, Count * sizeof(int), stream);
-  device->api->copyFromAsync(target, buffer, Count * sizeof(int), stream);
+  void* stream = device->api().createStream();
+  device->api().copyToAsync(buffer, source, Count * sizeof(int), stream);
+  device->api().copyFromAsync(target, buffer, Count * sizeof(int), stream);
 
   int seen = 0;
-  device->api->streamHostFunction(stream, [&]() { seen = target[Count - 1]; });
-  device->api->syncStreamWithHost(stream);
+  device->api().streamHostFunction(stream, [&]() { seen = target[Count - 1]; });
+  device->api().syncStreamWithHost(stream);
 
   EXPECT_EQ(1904, seen);
 
-  device->api->destroyGenericStream(stream);
-  device->api->freeGlobMem(buffer);
-  device->api->freePinnedMem(target);
-  device->api->freePinnedMem(source);
+  device->api().destroyGenericStream(stream);
+  device->api().freeGlobMem(buffer);
+  device->api().freePinnedMem(target);
+  device->api().freePinnedMem(source);
 }
 
 TEST_F(Streams, streamWorkIsDoneAfterSync) {
   constexpr std::size_t Count = 1024;
 
-  auto* source = static_cast<int*>(device->api->allocPinnedMem(Count * sizeof(int)));
-  auto* buffer = static_cast<int*>(device->api->allocGlobMem(Count * sizeof(int)));
+  auto* source = static_cast<int*>(device->api().allocPinnedMem(Count * sizeof(int)));
+  auto* buffer = static_cast<int*>(device->api().allocGlobMem(Count * sizeof(int)));
   std::fill(source, source + Count, 1904);
 
-  void* stream = device->api->createStream();
-  device->api->copyToAsync(buffer, source, Count * sizeof(int), stream);
-  device->api->syncStreamWithHost(stream);
+  void* stream = device->api().createStream();
+  device->api().copyToAsync(buffer, source, Count * sizeof(int), stream);
+  device->api().syncStreamWithHost(stream);
 
   // a bounded number of polls, so that a stream that never reports as done fails the test instead
   // of hanging it
   bool done = false;
   for (int poll = 0; poll < 1000 && !done; ++poll) {
-    done = device->api->isStreamWorkDone(stream);
+    done = device->api().isStreamWorkDone(stream);
   }
   EXPECT_TRUE(done);
 
-  device->api->destroyGenericStream(stream);
-  device->api->freeGlobMem(buffer);
-  device->api->freePinnedMem(source);
+  device->api().destroyGenericStream(stream);
+  device->api().freeGlobMem(buffer);
+  device->api().freePinnedMem(source);
 }
 
 TEST_F(Streams, destroyedStreamIsNotSynchronized) {
-  void* stream = device->api->createStream();
-  device->api->syncStreamWithHost(stream);
-  device->api->destroyGenericStream(stream);
+  void* stream = device->api().createStream();
+  device->api().syncStreamWithHost(stream);
+  device->api().destroyGenericStream(stream);
 
   // synchronizes all streams the device knows of; a destroyed stream must not be among them
-  device->api->syncDevice();
+  device->api().syncDevice();
   SUCCEED();
 }