From 1bfe88164693b154cd67365eed982b95da42547a Mon Sep 17 00:00:00 2001 From: David Schneller Date: Fri, 25 Sep 2026 20:39:50 +0000 Subject: [PATCH 1/6] refactor: share the explicit instantiations of the algorithms Every backend instantiates the member templates of device::Algorithms explicitly, and each of them spelled out its own list of 36 instantiations. algorithms/Instantiations.h now holds the type lists together with one instantiation macro per member template; the CUDA/HIP and the SYCL backend use them. A backend that is added later thus exports the same symbols by construction. No functional change: compiled against Algorithms.h, the explicit instantiations and the macros emit the same 36 symbols for both backends. AI-generated. Model: Claude Opus 5.5 --- README.md | 4 +- algorithms/Instantiations.h | 65 ++++++++++++++++++++++ algorithms/cudahip/ArrayManip.cpp | 33 +---------- algorithms/cudahip/BatchManip.cpp | 31 +---------- algorithms/cudahip/Debugging.cpp | 10 +--- algorithms/cudahip/Reduction.cpp | 92 +------------------------------ algorithms/sycl/ArrayManip.cpp | 33 +---------- algorithms/sycl/BatchManip.cpp | 31 +---------- algorithms/sycl/Debugging.cpp | 10 +--- algorithms/sycl/Reduction.cpp | 92 +------------------------------ 10 files changed, 88 insertions(+), 313 deletions(-) create mode 100644 algorithms/Instantiations.h diff --git a/README.md b/README.md index 7832ef8..7dd3d14 100644 --- a/README.md +++ b/README.md @@ -98,7 +98,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/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, size_t, void*); + +#define DEVICE_ALGORITHMS_INSTANTIATE_FILL_ARRAY(T) \ + template void Algorithms::fillArray(T*, T, size_t, void*); + +#define DEVICE_ALGORITHMS_INSTANTIATE_SET_TO_VALUE(T) \ + template void Algorithms::setToValue(T**, T, size_t, size_t, void*); + +#define DEVICE_ALGORITHMS_INSTANTIATE_ACCUMULATE_BATCHED_DATA(T) \ + template void Algorithms::accumulateBatchedData(const T**, T**, size_t, size_t, void*); + +#define DEVICE_ALGORITHMS_INSTANTIATE_COMPARE_DATA_WITH_HOST(T) \ + template void Algorithms::compareDataWithHost(const T*, const T*, size_t, const std::string&); + +#define DEVICE_ALGORITHMS_INSTANTIATE_REDUCE_VECTOR(AccT, VecT) \ + template void Algorithms::reduceVector( \ + 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 #include @@ -28,22 +29,7 @@ void Algorithms::scaleArray(T* devArray, T scalar, const size_t numElements, voi kernel_scaleArray<<>>(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 @@ -63,20 +49,7 @@ void Algorithms::fillArray(T* devArray, const T scalar, const size_t numElements kernel_fillArray<<>>(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 #include @@ -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 @@ -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 #include @@ -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/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 @@ -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 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 @@ -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 From e00cf240e20e6e0a79cf8ac3110bf6c8a9d745ee Mon Sep 17 00:00:00 2001 From: David Schneller Date: Fri, 25 Sep 2026 20:42:06 +0000 Subject: [PATCH 2/6] refactor!: access the device instance through accessors DeviceInstance::instance() returns the singleton; api() and algorithms() return references to the API and the algorithms of the instance. instance() is defined in device.cpp. All shared objects of a process thus see the same instance, also those compiled with hidden visibility, as Python modules built by pybind11 are. api() and algorithms() are const members that still give mutable access, as a pointer member does; call sites that hold a const reference to the instance keep working. This changes the interface for all consumers: getInstance() becomes instance(), api-> becomes api(). and algorithms. becomes algorithms(). The tests and the examples are migrated accordingly. AI-generated. Model: Claude Opus 5.5 --- UsmAllocator.h | 2 +- device.cpp | 27 +-- device.h | 20 ++- examples/basic/basic.cpp | 36 ++-- examples/jacobi/src/gpu/solver.cpp | 74 ++++---- examples/jacobi/src/helper.hpp | 4 +- examples/jacobi/tests/driver.cpp | 6 +- .../jacobi/tests/gpu/subroutines_tests.cpp | 4 +- examples/mpi/src/main.cpp | 8 +- tests/BaseTestSuite.h | 2 +- tests/array_manip.cpp | 78 +++++---- tests/batch_manip.cpp | 158 +++++++++--------- tests/main.cpp | 10 +- tests/memory.cpp | 24 +-- tests/reductions.cpp | 56 ++++--- tests/streams.cpp | 50 +++--- 16 files changed, 291 insertions(+), 268 deletions(-) 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/device.cpp b/device.cpp index f1b5182..9022cc4 100644 --- a/device.cpp +++ b/device.cpp @@ -14,20 +14,23 @@ #error "Unknown interface for the device wrapper" #endif +#include + 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()), algorithmsP(std::make_unique()) { + 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 + 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 apiP; + std::unique_ptr algorithmsP; }; } // namespace device diff --git a/examples/basic/basic.cpp b/examples/basic/basic.cpp index ac9d5ea..be44c36 100644 --- a/examples/basic/basic.cpp +++ b/examples/basic/basic.cpp @@ -15,44 +15,44 @@ int main(int argc, char* argv[]) { real* inputArray = new real[size]; real* outputArray = new real[size]; - 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; + std::cout << "max shared mem: " << device.api().getMaxSharedMemSize() << std::endl; + std::cout << "max thread block size: " << device.api().getMaxThreadBlockSize() << std::endl; // allocate mem. on a device - real* dInputArray = static_cast(device.api->allocGlobMem(sizeof(real) * size)); - real* dOutputArray = static_cast(device.api->allocGlobMem(sizeof(real) * size)); + real* dInputArray = static_cast(device.api().allocGlobMem(sizeof(real) * size)); + real* dOutputArray = static_cast(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(); + device.api().checkOffloading(); + device.api().syncDevice(); // copy data from a device - device.api->copyFrom(outputArray, dOutputArray, sizeof(real) * size); + device.api().copyFrom(outputArray, dOutputArray, sizeof(real) * size); // 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(); 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(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(devX, devTempX); - device.api->copyFrom(const_cast(x.data()), devTempX, x.size() * sizeof(real)); + device.api().copyFrom(const_cast(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(device.api->allocGlobMem(rhs.size() * sizeof(real))); - device.api->copyTo(devRhs, rhs.data(), rhs.size() * sizeof(real)); + devRhs = static_cast(device.api().allocGlobMem(rhs.size() * sizeof(real))); + device.api().copyTo(devRhs, rhs.data(), rhs.size() * sizeof(real)); - devX = static_cast(device.api->allocGlobMem(x.size() * sizeof(real))); - device.api->copyTo(devX, x.data(), x.size() * sizeof(real)); + devX = static_cast(device.api().allocGlobMem(x.size() * sizeof(real))); + device.api().copyTo(devX, x.data(), x.size() * sizeof(real)); - devTempX = static_cast(device.api->allocGlobMem(x.size() * sizeof(real))); - device.api->copyTo(devTempX, x.data(), x.size() * sizeof(real)); + devTempX = static_cast(device.api().allocGlobMem(x.size() * sizeof(real))); + device.api().copyTo(devTempX, x.data(), x.size() * sizeof(real)); - devTemp = static_cast(device.api->allocGlobMem(x.size() * sizeof(real))); + devTemp = static_cast(device.api().allocGlobMem(x.size() * sizeof(real))); // InvDiag still holds the diagonal elements at this point - devDiag = static_cast(device.api->allocGlobMem(invDiag.size() * sizeof(real))); - device.api->copyTo(devDiag, invDiag.data(), invDiag.size() * sizeof(real)); + devDiag = static_cast(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(device.api->allocGlobMem(invDiag.size() * sizeof(real))); - device.api->copyTo(devInvDiag, invDiag.data(), invDiag.size() * sizeof(real)); + devInvDiag = static_cast(device.api().allocGlobMem(invDiag.size() * sizeof(real))); + device.api().copyTo(devInvDiag, invDiag.data(), invDiag.size() * sizeof(real)); - devResidual = static_cast(device.api->allocGlobMem(residual.size() * sizeof(real))); - device.api->copyTo(devResidual, residual.data(), residual.size() * sizeof(real)); + devResidual = static_cast(device.api().allocGlobMem(residual.size() * sizeof(real))); + device.api().copyTo(devResidual, residual.data(), residual.size() * sizeof(real)); devLU = std::make_unique(lu.info); // GpuMatrixDataT devLU(LU.Info); - devLU->data = static_cast(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(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(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(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/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/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 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 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 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 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 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 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 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(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(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 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(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 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(N, M, true, [&](float** batch, float* data) { std::vector 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(N, M, false, [&](float** batch, float* data) { std::vector 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(N, M, false, [&](float** batch, float* data) { std::vector hostVector(N * M, 1); - device->api->copyToAsync( - data, &hostVector[0], N * M * sizeof(float), device->api->getDefaultStream()); - device->algorithms.copyScatterToUniform( - const_cast(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(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(N, M, false, [&](float** batch, float* data) { std::vector 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(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 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(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 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(N, M, true, [&](long** batch, long* data) { std::vector 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(N, M, false, [&](long** batch, long* data) { std::vector 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(N, M, false, [&](long** batch, long* data) { std::vector hostVector(N * M, 1); - device->api->copyToAsync( - data, &hostVector[0], N * M * sizeof(long), device->api->getDefaultStream()); - device->algorithms.copyScatterToUniform( - const_cast(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(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(N, M, false, [&](long** batch, long* data) { std::vector 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/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(device->api->allocGlobMem(sizeof(unsigned) * size)); - device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size); + auto* devVector = + reinterpret_cast(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* testResult = reinterpret_cast(device->api->allocPinnedMem(sizeof(unsigned))); + unsigned* testResult = + reinterpret_cast(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 vector(size, 0); - auto* devVector = reinterpret_cast(device->api->allocGlobMem(sizeof(unsigned) * size)); - device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size); + auto* devVector = + reinterpret_cast(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::min(); auto expectedResult = std::accumulate(vector.begin(), vector.end(), initValue, max); - unsigned* testResult = reinterpret_cast(device->api->allocPinnedMem(sizeof(unsigned))); + unsigned* testResult = + reinterpret_cast(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(device->api->allocGlobMem(sizeof(unsigned) * size)); - device->api->copyTo(devVector, vector.data(), sizeof(unsigned) * size); + auto* devVector = + reinterpret_cast(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::max(); auto expectedResult = std::accumulate(vector.begin(), vector.end(), initValue, min); - unsigned* testResult = reinterpret_cast(device->api->allocPinnedMem(sizeof(unsigned))); + unsigned* testResult = + reinterpret_cast(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(device->api->allocPinnedMem(Count * sizeof(int))); - auto* target = static_cast(device->api->allocPinnedMem(Count * sizeof(int))); - auto* buffer = static_cast(device->api->allocGlobMem(Count * sizeof(int))); + auto* source = static_cast(device->api().allocPinnedMem(Count * sizeof(int))); + auto* target = static_cast(device->api().allocPinnedMem(Count * sizeof(int))); + auto* buffer = static_cast(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(device->api->allocPinnedMem(Count * sizeof(int))); - auto* buffer = static_cast(device->api->allocGlobMem(Count * sizeof(int))); + auto* source = static_cast(device->api().allocPinnedMem(Count * sizeof(int))); + auto* buffer = static_cast(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(); } From 2f81a508bf1fde11b959f90516554f9f2e65551f Mon Sep 17 00:00:00 2001 From: David Schneller Date: Fri, 25 Sep 2026 21:03:14 +0000 Subject: [PATCH 3/6] feat: add a host backend, built with DEVICE_BACKEND=none With DEVICE_BACKEND=none, the device library implements its API on the host; code that uses the API thus builds and runs without any GPU toolchain or GPU. DEVICE_ARCH is not needed for it. The host backend has exactly one device, the host itself. Device memory is host memory, aligned to 128 bytes as on the GPU backends, with the same bookkeeping. All work runs synchronously on the calling thread: streams and events only serve as handles, host functions run right away, and streamWaitMemory makes the caller wait. Graph capturing is not available. The algorithms are plain loops with the semantics of the GPU kernels. The complete test suite passes on it, also under AddressSanitizer and UndefinedBehaviorSanitizer, and the library only depends on the C and C++ runtime libraries. AI-generated. Model: Claude Opus 5.5 --- CMakeLists.txt | 12 ++- README.md | 13 ++- algorithms/host/ArrayManip.cpp | 44 +++++++++ algorithms/host/BatchManip.cpp | 84 ++++++++++++++++++ algorithms/host/Debugging.cpp | 47 ++++++++++ algorithms/host/Reduction.cpp | 60 +++++++++++++ device.cpp | 2 + host.cmake | 14 +++ interfaces/host/Control.cpp | 107 ++++++++++++++++++++++ interfaces/host/HostWrappedAPI.h | 147 +++++++++++++++++++++++++++++++ interfaces/host/Memory.cpp | 146 ++++++++++++++++++++++++++++++ interfaces/host/Streams.cpp | 96 ++++++++++++++++++++ tests/CMakeLists.txt | 2 +- 13 files changed, 768 insertions(+), 6 deletions(-) create mode 100644 algorithms/host/ArrayManip.cpp create mode 100644 algorithms/host/BatchManip.cpp create mode 100644 algorithms/host/Debugging.cpp create mode 100644 algorithms/host/Reduction.cpp create mode 100644 host.cmake create mode 100644 interfaces/host/Control.cpp create mode 100644 interfaces/host/HostWrappedAPI.h create mode 100644 interfaces/host/Memory.cpp create mode 100644 interfaces/host/Streams.cpp 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 7dd3d14..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 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 +#include +#include +#include + +namespace device { +template +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(devArray[i] * scalar); + } +} +DEVICE_ALGORITHMS_ARRAY_TYPES(DEVICE_ALGORITHMS_INSTANTIATE_SCALE_ARRAY) + +template +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(base); + for (size_t i = 0; i < numElements; ++i) { + out[i] = reinterpret_cast(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 +#include +#include + +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 +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 +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(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(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 +#include +#include +#include + +namespace device { + +template +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 +#include + +namespace device { +namespace { +template +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(buffer[i])); + } + *result = accumulator; +} +} // namespace + +template +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::lowest(), + [](AccT a, AccT b) { return a > b ? a : b; }); + break; + } + case ReductionType::Min: { + reduce( + result, buffer, overrideResult, size, std::numeric_limits::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/device.cpp b/device.cpp index 9022cc4..51851a6 100644 --- a/device.cpp +++ b/device.cpp @@ -10,6 +10,8 @@ #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 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 +#include +#include +#include + +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(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(pages) * static_cast(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 +#include +#include +#include +#include +#include +#include +#include + +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& 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& 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> genericStreams{}; + + // the same bookkeeping as in the other backends + size_t allocatedMemBytes{0}; + size_t allocatedUnifiedMemBytes{0}; + size_t deallocatedMemBytes{0}; + std::unordered_map 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 +#include +#include +#include +#include + +using namespace device; + +void* ConcreteAPI::allocate(size_t size, bool unified) { + isFlagSet(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(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(status); + std::lock_guard guard(apiMutex); + std::ostringstream report{}; + report << "Memory Leaks, bytes: " << (allocatedMemBytes - deallocatedMemBytes) << '\n'; + return report.str(); +} + +size_t ConcreteAPI::getCurrentlyOccupiedMem() { + isFlagSet(status); + std::lock_guard guard(apiMutex); + return allocatedMemBytes; +} + +size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { + isFlagSet(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(dst); + const auto* srcBytes = static_cast(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 +#include +#include +#include +#include +#include +#include + +using namespace device; + +ConcreteAPI::Event* ConcreteAPI::toEvent(void* eventPtr) { return static_cast(eventPtr); } + +void* ConcreteAPI::getDefaultStream() { + isFlagSet(status); + return &defaultStream; +} + +void ConcreteAPI::syncDefaultStreamWithHost() { isFlagSet(status); } + +void* ConcreteAPI::createStream(double /*priority*/) { + isFlagSet(status); + auto stream = std::make_unique(); + void* streamPtr = stream.get(); + std::lock_guard guard(apiMutex); + genericStreams.emplace(streamPtr, std::move(stream)); + return streamPtr; +} + +void ConcreteAPI::destroyGenericStream(void* streamPtr) { + isFlagSet(status); + std::lock_guard guard(apiMutex); + genericStreams.erase(streamPtr); +} + +void ConcreteAPI::syncStreamWithHost(void* /*streamPtr*/) { + isFlagSet(status); +} + +bool ConcreteAPI::isStreamWorkDone(void* /*streamPtr*/) { + isFlagSet(status); + return true; +} + +void ConcreteAPI::syncStreamWithEvent(void* /*streamPtr*/, void* /*eventPtr*/) {} + +void ConcreteAPI::streamHostFunction(void* /*streamPtr*/, const std::function& 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& /*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(span).count(); +} diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index ece56b8..06060da 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) From ee2e0207de8ddef0357895598b84f0f455fbd7c9 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Fri, 25 Sep 2026 21:05:30 +0000 Subject: [PATCH 4/6] test: check the guarantees of the host backend The common test suite does not cover events, streamWaitMemory, the order of several host functions on one stream, the alignment and the bookkeeping of allocations, or copies with a pitch that differs from the width. tests/host.cpp checks these for the host backend, together with its properties and the neutral elements of the reductions; it is only built with DEVICE_BACKEND=none. Both a reduction that starts the maximum at the smallest positive value and a 2D copy that ignores the pitches make these tests fail, while the common suite still passes. AI-generated. Model: Claude Opus 5.5 --- tests/CMakeLists.txt | 7 ++ tests/host.cpp | 193 +++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 200 insertions(+) create mode 100644 tests/host.cpp diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 06060da..7c7b99a 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -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/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 +#include +#include +#include +#include +#include + +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(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 order; + api.streamHostFunction(stream, [&] { order.push_back(1); }); + api.streamHostFunction(stream, [&] { order.push_back(2); }); + api.syncStreamWithHost(stream); + EXPECT_EQ(order, (std::vector{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( + 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 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 values{-3.0F, -1.5F, -7.25F}; + const auto bytes = values.size() * sizeof(float); + auto* buffer = static_cast(api.allocGlobMem(bytes)); + api.copyTo(buffer, values.data(), bytes); + auto* result = static_cast(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 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('a' + row * Width + column); + } + } + + auto* dev = static_cast(api.allocGlobMem(DevPitch * Height)); + std::memset(dev, Padding, DevPitch * Height); + api.copy2dArrayTo(dev, DevPitch, src.data(), SrcPitch, Width, Height); + + std::vector 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); +} From bfee83336cae670187bb1d3ef34b36b1cf06da7e Mon Sep 17 00:00:00 2001 From: David Schneller Date: Fri, 25 Sep 2026 21:06:38 +0000 Subject: [PATCH 5/6] ci: build and test the host backend Adds DEVICE_BACKEND=none to the build matrix. The job needs no GPU; it runs on a GitHub-hosted runner in the CPU image that SeisSol uses as well, and it runs the complete test suite, including the tests of the host backend. AI-generated. Model: Claude Opus 5.5 --- .github/workflows/test.yml | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/.github/workflows/test.yml b/.github/workflows/test.yml index bdd29f5..a7f9024 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 From 857b9f6947009b90dfbe329610e2ebf5b9ba3edc Mon Sep 17 00:00:00 2001 From: David Schneller Date: Fri, 25 Sep 2026 21:07:41 +0000 Subject: [PATCH 6/6] fix: make the basic example build and check itself The basic example used the real type without including common.h, and it called getMaxSharedMemSize, getMaxThreadBlockSize and checkOffloading, none of which the API provides. It now copies initialized data to the device, within the device and back, and fails if the result differs. The CI job of the host backend builds and runs it. AI-generated. Model: Claude Opus 5.5 --- .github/workflows/test.yml | 17 +++++++++++++++++ examples/basic/basic.cpp | 24 ++++++++++++++++++------ examples/cmake/common_options.cmake | 2 +- 3 files changed, 36 insertions(+), 7 deletions(-) diff --git a/.github/workflows/test.yml b/.github/workflows/test.yml index a7f9024..133a8e7 100644 --- a/.github/workflows/test.yml +++ b/.github/workflows/test.yml @@ -157,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/examples/basic/basic.cpp b/examples/basic/basic.cpp index be44c36..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 #include -#include using namespace device; @@ -14,6 +14,10 @@ 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(i); + outputArray[i] = 0; + } DeviceInstance& device = DeviceInstance::instance(); @@ -31,8 +35,6 @@ int main(int argc, char* argv[]) { 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; // allocate mem. on a device real* dInputArray = static_cast(device.api().allocGlobMem(sizeof(real) * size)); @@ -41,13 +43,21 @@ int main(int argc, char* argv[]) { // copy data into a device device.api().copyTo(dInputArray, inputArray, sizeof(real) * size); - // call a kernel - device.api().checkOffloading(); + // 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); + 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); @@ -58,4 +68,6 @@ int main(int argc, char* argv[]) { 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"