diff --git a/.github/workflows/test.yml b/.github/workflows/test.yml index bdd29f5..2bafc70 100644 --- a/.github/workflows/test.yml +++ b/.github/workflows/test.yml @@ -35,6 +35,7 @@ jobs: container-options: --runtime=nvidia --gpus=all pythonbreak: true test: true + sanitizer: compute-sanitizer - arch: sm_86 backend: acpp cc: gcc-13 @@ -55,6 +56,7 @@ jobs: container-options: --runtime=nvidia --gpus=all pythonbreak: true test: true + sanitizer: compute-sanitizer - arch: sm_86 backend: cuda cc: nvc @@ -122,8 +124,7 @@ jobs: cd tests mkdir build && cd build - # TODO: change to `-Wall -Werror` at some point - EXTRA_FLAGS="-Wall" + EXTRA_FLAGS="-Wall -Werror" export CFLAGS="${EXTRA_FLAGS} ${CFLAGS}" export CXXFLAGS="${EXTRA_FLAGS} ${CXXFLAGS}" @@ -148,3 +149,15 @@ jobs: cd build ./tests + + - id: sanitize + name: sanitize-device + if: ${{matrix.setup.test && matrix.setup.sanitizer && matrix.build_type == 'Debug'}} + run: | + cd tests + cd build + + # The three large reductions take minutes under the sanitizer and cover no pointer + # arithmetic the smaller suites do not, so they sit this one out. + ${{matrix.setup.sanitizer}} --tool memcheck --error-exitcode 1 ./tests \ + --gtest_filter=-Reductions.Add:Reductions.Max:Reductions.Min diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml deleted file mode 100644 index feb69cd..0000000 --- a/.gitlab-ci.yml +++ /dev/null @@ -1,89 +0,0 @@ -# SPDX-FileCopyrightText: 2019 SeisSol Group -# -# SPDX-License-Identifier: BSD-3-Clause - -stages: - - adjust - - test - - miscellaneous - - -image: - name: "seissol/gpu-ci-image:v0.0.4" - entrypoint: [""] - - -submodules: - stage: adjust - tags: - - atsccs68-docker-executor - script: - - cat ./.gitmodules - - git submodule update --init --recursive - artifacts: - paths: - - submodules - - -.common_test_steps: &common_test_script - - export CTEST_OUTPUT_ON_FAILURE=1 - - if [[ "${BACKEND}" == "hip" ]]; then - . /etc/profile.d/rocm.sh ; - export HIP_PLATFORM=nvidia ; - fi ; - - -tests: - stage: test - tags: - - atsccs68-docker-executor - parallel: - matrix: - - BACKEND: [cuda, hip, hipsycl] - script: - - *common_test_script - - mkdir -p tests/build - - cd tests/build - - cmake .. -DDEVICE_BACKEND=${BACKEND} -DSM=${GPU_MODEL} -DREAL_SIZE_IN_BYTES=4 - - make -j4 - - ./tests - - -basic_example: - stage: miscellaneous - tags: - - atsccs68-docker-executor - parallel: - matrix: - - BACKEND: [cuda, hip, hipsycl] - script: - - *common_test_script - - export WORKDIR=./examples/basic/build - - mkdir -p ${WORKDIR} && cd ${WORKDIR} - - cmake .. -DDEVICE_BACKEND=${BACKEND} -DSM=${GPU_MODEL} -DREAL_SIZE_IN_BYTES=4 - - make -j4 - - ./basic - - -jacobi_example: - stage: miscellaneous - tags: - - atsccs68-docker-executor - parallel: - matrix: - - BACKEND: [cuda, hip, hipsycl] - before_script: - - git clone --depth 1 --branch release-1.10.0 https://github.com/google/googletest - - mkdir -p googletest/build && cd googletest/build - - cmake .. -DBUILD_GTEST=ON -DBUILD_GMOCK=ON -Dgtest_disable_pthreads=ON -DBUILD_SHARED_LIBS=ON - - make -j $(nproc) - - make install - - cd ../.. - script: - - *common_test_script - - export WORKDIR=./examples/jacobi/build - - mkdir -p ${WORKDIR} && cd ${WORKDIR} - - cmake .. -DDEVICE_BACKEND=${BACKEND} -DSM=${GPU_MODEL} -DREAL_SIZE_IN_BYTES=4 -DWITH_MPI=OFF -DWITH_TESTS=ON - - make -j4 - - ./tests - - ./solver ./config.yaml diff --git a/AbstractAPI.h b/AbstractAPI.h index bde98ee..f5b5178 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -33,6 +33,11 @@ enum class ProfilingColors : uint32_t { struct AbstractAPI { virtual ~AbstractAPI() = default; + /** + * Selects the device for the calling thread and makes it the choice of the process. Backends + * that keep the selected device per thread give a thread that has not called this the device + * the process selected, the first time that thread asks for the device id. + */ virtual void setDevice(int deviceId) = 0; virtual int getDeviceId() = 0; @@ -85,10 +90,59 @@ struct AbstractAPI { virtual void syncDefaultStreamWithHost() = 0; virtual bool isCapableOfGraphCapturing() = 0; - virtual DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) = 0; - virtual void streamEndCapture(DeviceGraphHandle handle) = 0; - virtual void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) = 0; - + virtual DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) = 0; + virtual void streamEndCapture(const DeviceGraphHandle& handle) = 0; + virtual void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) = 0; + + /** + * Explicit graph construction. + * + * Instead of recording a whole stream and letting the backend infer the dependency structure + * from events, the caller states the structure directly: every graphAddNode call contributes + * the work recorded by `recorder` and makes it depend on exactly `dependencies`. Fork/join is + * then a property of the graph rather than something that has to be expressed through streams + * and events. + * + * A single graph is built by one thread at a time. `recorder` receives a stream that is only a + * recording vehicle: what it enqueues becomes the node, and the dependencies stated in the call + * are the ones the node is guaranteed to get. + * + * Nodes that are meant to run concurrently have to be recorded onto different streams. A stream + * can be reused for a later node, but a backend that expresses edges through the recorded + * stream rather than through node handles - the SYCL one does, since the graph extension has no + * node handles to hand out - adds an edge between two nodes that shared a stream, and those two + * then run one after the other. + * + * If `recorder` enqueues nothing, the returned handle refers to `dependencies` themselves, so + * an empty recorder is a valid way to express a pure join node. + * + * graphBeginNode and graphEndNode are the same thing split in two, for callers that cannot + * wrap the recorded work in a callback and instead have to leave a node open across code they + * do not control. Only one node per stream may be open at a time. + */ + virtual bool isCapableOfGraphNodes() = 0; + virtual DeviceGraphHandle graphCreate() = 0; + virtual void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) = 0; + virtual DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) = 0; + virtual void graphInstantiate(const DeviceGraphHandle& graphHandle) = 0; + + DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) { + graphBeginNode(graphHandle, dependencies, streamPtr); + recorder(streamPtr); + return graphEndNode(graphHandle, streamPtr); + } + + /** + * Creates a stream. `priority` runs from 0 for the lowest to 1 for the highest priority the + * device offers; NAN asks for the runtime default. Backends that only know a few priority + * classes round to the nearest one, and a device without priority support ignores the value. + */ virtual void* createStream(double priority = NAN) = 0; virtual void destroyGenericStream(void* streamPtr) = 0; virtual void syncStreamWithHost(void* streamPtr) = 0; @@ -96,6 +150,11 @@ struct AbstractAPI { virtual void syncStreamWithEvent(void* streamPtr, void* eventPtr) = 0; virtual void streamHostFunction(void* streamPtr, const std::function& function) = 0; + /** + * Blocks the stream until the value at `location` has reached at least `value`. `location` has + * to be host memory that the device can read, i.e. an allocation from allocPinnedMem with + * Destination::CurrentDevice. + */ virtual void streamWaitMemory(void* streamPtr, uint32_t* location, uint32_t value) = 0; virtual void* createEvent(bool withTiming = false) = 0; diff --git a/CMakeLists.txt b/CMakeLists.txt index e58157a..24c1f69 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -53,8 +53,8 @@ elseif((${DEVICE_BACKEND} STREQUAL "oneapi") OR (${DEVICE_BACKEND} STREQUAL "hip include(sycl.cmake) endif() -# common options -target_compile_features(device PRIVATE cxx_std_17) +# common options; the public headers use C++17 as well, so consumers get it too +target_compile_features(device PUBLIC cxx_std_17) if (USE_GRAPH_CAPTURING) target_compile_definitions(device PRIVATE DEVICE_USE_GRAPH_CAPTURING) diff --git a/DataTypes.h b/DataTypes.h index 9f707cb..5aa723f 100644 --- a/DataTypes.h +++ b/DataTypes.h @@ -7,28 +7,64 @@ #include #include +#include namespace device { -struct DeviceGraphHandle { - static const size_t invalidId{std::numeric_limits::max()}; - public: - explicit DeviceGraphHandle() : graphId(invalidId) {} - explicit DeviceGraphHandle(size_t id) : graphId(id) {} +/** + * Backend-specific payload of a compute graph. Only the active interface implementation defines + * this type; every other translation unit sees an incomplete type and reaches the graph through + * DeviceGraphHandle. + */ +struct DeviceGraph; - DeviceGraphHandle(const DeviceGraphHandle& other) = default; - DeviceGraphHandle& operator=(const DeviceGraphHandle& other) = default; +/** + * Owning handle to a compute graph. + * + * The backend resources (the graph and its executable instance) are released once the last handle + * pointing to them goes out of scope. A graph that is dropped from a cache therefore also frees + * its device-side resources. + */ +class DeviceGraphHandle { + public: + DeviceGraphHandle() = default; + explicit DeviceGraphHandle(std::shared_ptr graphPtr) : graph(std::move(graphPtr)) {} - bool isInitialized() const { return graphId != invalidId; } + [[nodiscard]] bool isInitialized() const { return static_cast(graph); } operator bool() const { return isInitialized(); } bool operator!() const { return !isInitialized(); } - size_t getGraphId() { return graphId; } + [[nodiscard]] DeviceGraph* get() const { return graph.get(); } + + void reset() { graph.reset(); } + + private: + std::shared_ptr graph; +}; + +/** + * Refers to the set of graph nodes produced by a single AbstractAPI::graphAddNode call. + * + * A node handle is an index into the graph that produced it and stays valid for that graph's + * lifetime. Passing it to a different graph is undefined. + */ +class DeviceGraphNodeHandle { + public: + static const size_t invalidId{std::numeric_limits::max()}; + + DeviceGraphNodeHandle() = default; + explicit DeviceGraphNodeHandle(size_t id) : nodeId(id) {} + + [[nodiscard]] bool isInitialized() const { return nodeId != invalidId; } + + operator bool() const { return isInitialized(); } + + [[nodiscard]] size_t getNodeId() const { return nodeId; } private: - size_t graphId{invalidId}; + size_t nodeId{invalidId}; }; } // namespace device diff --git a/UsmAllocator.h b/UsmAllocator.h index abb9efb..37d9ba0 100644 --- a/UsmAllocator.h +++ b/UsmAllocator.h @@ -44,7 +44,7 @@ class UsmAllocator { void deallocate(T* ptr, std::size_t) { if (ptr) { - api->freeGlobMem(ptr); + api->freeUnifiedMem(ptr); } } diff --git a/algorithms/Common.h b/algorithms/Common.h index 762b4f0..b553cf6 100644 --- a/algorithms/Common.h +++ b/algorithms/Common.h @@ -9,7 +9,10 @@ #include "Algorithms.h" #include "Internals.h" +#include +#include #include +#include #if defined(__ACPP__) #include @@ -57,14 +60,21 @@ DEVICE_DEVICEFUNC T ntload(const T* location) { template int blockcount(F&& func, int blocksize) { - int device = 0; - int smCount = 0; - int blocksPerSM = 0; - APIWRAP(cudaGetDevice(&device)); - APIWRAP(cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, device)); - APIWRAP(cudaOccupancyMaxActiveBlocksPerMultiprocessor( - &blocksPerSM, std::forward(func), blocksize, 0)); - return smCount * blocksPerSM; + // The occupancy of a kernel does not change while the program runs, and this sits in front of + // every algorithm launch, so it is asked once per kernel. Every caller uses the default block + // size, and the devices of a node are the same model. + static const int count = [&]() { + int device = 0; + int smCount = 0; + int blocksPerSM = 0; + APIWRAP(cudaGetDevice(&device)); + APIWRAP(cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, device)); + APIWRAP(cudaOccupancyMaxActiveBlocksPerMultiprocessor( + &blocksPerSM, std::forward(func), blocksize, 0)); + // a grid of zero blocks is not a valid launch configuration + return std::max(1, smCount * blocksPerSM); + }(); + return count; } using LocalInt4 = int4; @@ -83,14 +93,21 @@ DEVICE_DEVICEFUNC T ntload(const T* location) { template int blockcount(F&& func, int blocksize) { - int device = 0; - int smCount = 0; - int blocksPerSM = 0; - APIWRAP(hipGetDevice(&device)); - APIWRAP(hipDeviceGetAttribute(&smCount, hipDeviceAttributeMultiprocessorCount, device)); - APIWRAP(hipOccupancyMaxActiveBlocksPerMultiprocessor( - &blocksPerSM, std::forward(func), blocksize, 0)); - return smCount * blocksPerSM; + // The occupancy of a kernel does not change while the program runs, and this sits in front of + // every algorithm launch, so it is asked once per kernel. Every caller uses the default block + // size, and the devices of a node are the same model. + static const int count = [&]() { + int device = 0; + int smCount = 0; + int blocksPerSM = 0; + APIWRAP(hipGetDevice(&device)); + APIWRAP(hipDeviceGetAttribute(&smCount, hipDeviceAttributeMultiprocessorCount, device)); + APIWRAP(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &blocksPerSM, std::forward(func), blocksize, 0)); + // a grid of zero blocks is not a valid launch configuration + return std::max(1, smCount * blocksPerSM); + }(); + return count; } using LocalInt4 = __attribute__((vector_size(16))) int; @@ -140,12 +157,35 @@ DEVICE_DEVICEFUNC std::size_t iimemcpy(void* dst, return end * sizeof(T); } +/** + * Returns the address bits that keep the given pointers from being 16-byte aligned; zero means + * both of them are. A null pointer contributes nothing, which is how the single-pointer case is + * asked. + * + * The copy and fill routines below step down from 16-byte accesses to single bytes. Batched + * buffers are addressed through a pointer table and an element stride, so their elements are not + * guaranteed to sit on a 16-byte boundary, and a vector access to an address that is not aligned + * to its own width faults. + */ +DEVICE_DEVICEFUNC std::size_t unalignedBits(const void* first, const void* second) { + const auto bits = + reinterpret_cast(first) | reinterpret_cast(second); + return static_cast(bits & 15U); +} + DEVICE_DEVICEFUNC void imemcpy(void* dst, const void* src, std::size_t count, int local, std::size_t stride) { + const auto lowBits = unalignedBits(dst, src); std::size_t offset = 0; - offset += iimemcpy(dst, src, offset, count, local, stride); - offset += iimemcpy(dst, src, offset, count, local, stride); - offset += iimemcpy(dst, src, offset, count, local, stride); + if (lowBits % sizeof(LocalInt4) == 0) { + offset += iimemcpy(dst, src, offset, count, local, stride); + } + if (lowBits % sizeof(LocalInt2) == 0) { + offset += iimemcpy(dst, src, offset, count, local, stride); + } + if (lowBits % sizeof(int) == 0) { + offset += iimemcpy(dst, src, offset, count, local, stride); + } offset += iimemcpy(dst, src, offset, count, local, stride); } @@ -165,10 +205,17 @@ DEVICE_DEVICEFUNC std::size_t } DEVICE_DEVICEFUNC void imemset(void* dst, std::size_t count, int local, std::size_t stride) { + const auto lowBits = unalignedBits(dst, nullptr); std::size_t offset = 0; - offset += iimemset(dst, offset, count, local, stride); - offset += iimemset(dst, offset, count, local, stride); - offset += iimemset(dst, offset, count, local, stride); + if (lowBits % sizeof(LocalInt4) == 0) { + offset += iimemset(dst, offset, count, local, stride); + } + if (lowBits % sizeof(LocalInt2) == 0) { + offset += iimemset(dst, offset, count, local, stride); + } + if (lowBits % sizeof(int) == 0) { + offset += iimemset(dst, offset, count, local, stride); + } offset += iimemset(dst, offset, count, local, stride); } diff --git a/algorithms/cudahip/BatchManip.cpp b/algorithms/cudahip/BatchManip.cpp index 9ce8b04..ce060ed 100644 --- a/algorithms/cudahip/BatchManip.cpp +++ b/algorithms/cudahip/BatchManip.cpp @@ -45,6 +45,9 @@ __global__ void kernel_accumulateBatchedData(const T** baseSrcPtr, for (size_t block = blockIdx.x; block < elementCount; block += gridDim.x) { const T* srcElement = baseSrcPtr[block]; T* dstElement = baseDstPtr[block]; + if (srcElement == nullptr || dstElement == nullptr) { + continue; + } #pragma unroll 4 for (int index = threadIdx.x; index < elementSize; index += device::internals::DefaultBlockDim) { @@ -105,6 +108,9 @@ template __global__ void kernel_setToValue(T** out, T value, size_t elementSize, size_t elementCount) { for (size_t block = blockIdx.x; block < elementCount; block += gridDim.x) { T* element = out[block]; + if (element == nullptr) { + continue; + } #pragma unroll 4 for (int index = threadIdx.x; index < elementSize; index += device::internals::DefaultBlockDim) { @@ -148,7 +154,9 @@ __global__ void kernel_copyUniformToScatter( const void* srcElement = reinterpret_cast(&reinterpret_cast(src)[block * srcOffset]); void* dstElement = dst[block]; - imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + if (dstElement != nullptr) { + imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + } } } @@ -172,7 +180,9 @@ __global__ void kernel_copyScatterToUniform( for (size_t block = blockIdx.x; block < elementCount; block += gridDim.x) { const void* srcElement = src[block]; void* dstElement = reinterpret_cast(&reinterpret_cast(dst)[block * dstOffset]); - imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + if (srcElement != nullptr) { + imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + } } } diff --git a/algorithms/cudahip/Reduction.cpp b/algorithms/cudahip/Reduction.cpp index 12f1bf7..40001af 100644 --- a/algorithms/cudahip/Reduction.cpp +++ b/algorithms/cudahip/Reduction.cpp @@ -7,7 +7,10 @@ #include #include +#include #include +#include +#include namespace device { @@ -21,7 +24,9 @@ struct Sum { template struct Max { - T defaultValue{std::numeric_limits::min()}; + // lowest(), not min(): for floating point types min() is the smallest positive normal value, + // which is larger than every negative input + T defaultValue{std::numeric_limits::lowest()}; __device__ __forceinline__ T operator()(T op1, T op2) { return op1 > op2 ? op1 : op2; } }; @@ -57,38 +62,79 @@ __device__ __forceinline__ T warpReduce(T value, OperationT operation) { return value; } -// Helper function for Generic Atomic Update -// Fallback to atomicCAS-based implementation if atomic instruction is not available -// Picked from: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#atomic-functions +template +struct AtomicWord { + static_assert(Size == 4 || Size == 8, "no atomic word of the size of the accumulator type"); +}; + +template <> +struct AtomicWord<4> { + using Type = unsigned int; +}; + +template <> +struct AtomicWord<8> { + using Type = unsigned long long; +}; + +template +__device__ __forceinline__ ToT reinterpretValue(const FromT& from) { + static_assert(sizeof(ToT) == sizeof(FromT), "reinterpretValue requires types of equal size"); + ToT to{}; + memcpy(&to, &from, sizeof(ToT)); + return to; +} + +// Fallback for the combinations without a native atomic. The compare-and-swap runs on a word of +// exactly the size of T: a wider word would read and write the memory next to the result. template -__device__ __forceinline__ void atomicUpdate(T* address, T val, OperationT operation) { - unsigned long long* address_as_ull = (unsigned long long*)address; - unsigned long long old = *address_as_ull, assumed; +__device__ __forceinline__ void atomicUpdateCas(T* address, T val, OperationT operation) { + using WordT = typename AtomicWord::Type; + auto* wordAddress = reinterpret_cast(address); + + WordT old = *wordAddress; + WordT assumed{}; do { assumed = old; - T calculatedRes = operation(*(T*)&assumed, val); - old = atomicCAS(address_as_ull, assumed, *(unsigned long long*)&calculatedRes); + const T updated = operation(reinterpretValue(assumed), val); + old = atomicCAS(wordAddress, assumed, reinterpretValue(updated)); } while (assumed != old); } -// Native atomics -template <> -__device__ __forceinline__ void - atomicUpdate>(int* address, int val, device::Sum operation) { - atomicAdd(address, val); -} -template <> -__device__ __forceinline__ void atomicUpdate>( - float* address, float val, device::Sum operation) { - atomicAdd(address, val); -} -#if __CUDA_ARCH__ >= 600 -template <> -__device__ __forceinline__ void atomicUpdate>( - double* address, double val, device::Sum operation) { - atomicAdd(address, val); -} +template +__device__ __forceinline__ void atomicUpdate(T* address, T val, OperationT operation) { + if constexpr (std::is_same_v>) { + if constexpr (std::is_same_v || std::is_same_v || + std::is_same_v || std::is_same_v) { + atomicAdd(address, val); + return; + } else if constexpr (std::is_integral_v && sizeof(T) == sizeof(unsigned long long)) { + // the unsigned addition wraps the same way, so it also gives the signed result + atomicAdd(reinterpret_cast(address), + static_cast(val)); + return; + } else if constexpr (std::is_same_v) { +// mirrors the guard the toolkit puts on the declaration itself +#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 600) + atomicAdd(address, val); + return; #endif + } + } + if constexpr (std::is_same_v> && + (std::is_same_v || std::is_same_v || + std::is_same_v)) { + atomicMax(address, val); + return; + } + if constexpr (std::is_same_v> && + (std::is_same_v || std::is_same_v || + std::is_same_v)) { + atomicMin(address, val); + return; + } + atomicUpdateCas(address, val, operation); +} // Block Reduce template @@ -122,8 +168,8 @@ __global__ void initKernel(T* result, OperationT operation) { } template -__launch_bounds__(BlockSize) void __global__ kernel_reduce( - AccT* result, const VecT* vector, size_t size, bool overrideResult, OperationT operation) { +__launch_bounds__(BlockSize) void __global__ + kernel_reduce(AccT* result, const VecT* vector, size_t size, OperationT operation) { // Maximum block size 1024, warp size 32 so 1024/32 = 32 chosen // For AMD, warp size 64, 1024/64 = 16, but 32 should work with a few idle memory addresses @@ -144,7 +190,6 @@ __launch_bounds__(BlockSize) void __global__ kernel_reduce( AccT blockAcc = blockReduce(threadAcc, shmem, operation); if (threadIdx.x == 0) { - (void)overrideResult; // to silence unused parameter warning for non-Add reductions atomicUpdate(result, blockAcc, operation); } } @@ -175,20 +220,24 @@ void Algorithms::reduceVector(AccT* result, } } + // the result is set either way, but there is nothing to reduce into it, and a grid of zero + // blocks is not a valid launch configuration + if (size == 0) { + CHECK_ERR; + return; + } + switch (type) { case ReductionType::Add: { - kernel_reduce<<>>( - result, buffer, size, overrideResult, device::Sum()); + kernel_reduce<<>>(result, buffer, size, device::Sum()); break; } case ReductionType::Max: { - kernel_reduce<<>>( - result, buffer, size, overrideResult, device::Max()); + kernel_reduce<<>>(result, buffer, size, device::Max()); break; } case ReductionType::Min: { - kernel_reduce<<>>( - result, buffer, size, overrideResult, device::Min()); + kernel_reduce<<>>(result, buffer, size, device::Min()); break; } default: { diff --git a/algorithms/sycl/ArrayManip.cpp b/algorithms/sycl/ArrayManip.cpp index 4a72d23..5b68249 100644 --- a/algorithms/sycl/ArrayManip.cpp +++ b/algorithms/sycl/ArrayManip.cpp @@ -14,6 +14,11 @@ using namespace device::internals; namespace device { template void Algorithms::scaleArray(T* devArray, T scalar, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, numElements); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -44,6 +49,11 @@ template void template void Algorithms::fillArray(T* devArray, const T scalar, const size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, numElements); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -72,6 +82,11 @@ template void Algorithms::fillArray(char* devArray, char scalar, const size_t numElements, void* streamPtr); void Algorithms::touchMemoryI(void* ptr, size_t size, bool clean, void* streamPtr) { + // an empty range is not a valid launch configuration + if (size == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, size); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -94,6 +109,11 @@ void Algorithms::incrementalAddI( uintptr_t* oout = reinterpret_cast(out); uintptr_t obase = reinterpret_cast(base); + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, numElements); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { diff --git a/algorithms/sycl/BatchManip.cpp b/algorithms/sycl/BatchManip.cpp index 0aab2b3..f129a47 100644 --- a/algorithms/sycl/BatchManip.cpp +++ b/algorithms/sycl/BatchManip.cpp @@ -17,6 +17,11 @@ void Algorithms::streamBatchedDataI(const void** baseSrcPtr, size_t elementSize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -39,6 +44,11 @@ void Algorithms::streamBatchedDataI(const void** baseSrcPtr, template void Algorithms::accumulateBatchedData( const T** baseSrcPtr, T** baseDstPtr, size_t elementSize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -69,6 +79,11 @@ template void Algorithms::accumulateBatchedData(const double** baseSrcPtr, void Algorithms::touchBatchedMemoryI( void** basePtr, size_t elementSize, size_t numElements, bool clean, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -93,6 +108,11 @@ void Algorithms::touchBatchedMemoryI( template void Algorithms::setToValue( T** out, T value, size_t elementSize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -134,6 +154,11 @@ void Algorithms::copyUniformToScatterI(const void* src, size_t copySize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -143,11 +168,13 @@ void Algorithms::copyUniformToScatterI(const void* src, const void* srcElement = reinterpret_cast(&reinterpret_cast(src)[block * srcOffset]); void* dstElement = dst[block]; - imemcpy(dstElement, - srcElement, - copySize, - item.get_local_id(0), - device::internals::DefaultBlockDim); + if (dstElement != nullptr) { + imemcpy(dstElement, + srcElement, + copySize, + item.get_local_id(0), + device::internals::DefaultBlockDim); + } }); }); } @@ -158,6 +185,11 @@ void Algorithms::copyScatterToUniformI(const void** src, size_t copySize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -166,11 +198,13 @@ void Algorithms::copyScatterToUniformI(const void** src, const auto block = item.get_group().get_group_id(0); const void* srcElement = src[block]; void* dstElement = reinterpret_cast(&reinterpret_cast(dst)[block * dstOffset]); - imemcpy(dstElement, - srcElement, - copySize, - item.get_local_id(0), - device::internals::DefaultBlockDim); + if (srcElement != nullptr) { + imemcpy(dstElement, + srcElement, + copySize, + item.get_local_id(0), + device::internals::DefaultBlockDim); + } }); }); } diff --git a/algorithms/sycl/Reduction.cpp b/algorithms/sycl/Reduction.cpp index e0d5d52..91421cd 100644 --- a/algorithms/sycl/Reduction.cpp +++ b/algorithms/sycl/Reduction.cpp @@ -10,8 +10,6 @@ #include #include -#if 1 - namespace { using namespace device; @@ -21,7 +19,9 @@ constexpr T neutral() { return T(0); } if constexpr (Type == ReductionType::Max) { - return std::numeric_limits::min(); + // lowest(), not min(): for floating point types min() is the smallest positive normal value, + // which is larger than every negative input + return std::numeric_limits::lowest(); } if constexpr (Type == ReductionType::Min) { return std::numeric_limits::max(); @@ -88,6 +88,11 @@ void launchReduction(AccT* result, }); } + // the result is set either way, but there is nothing to reduce into it + if (size == 0) { + return; + } + ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { const size_t numWorkGroups = (size + (workGroupSize * itemsPerWorkItem) - 1) / (workGroupSize * itemsPerWorkItem); @@ -148,58 +153,6 @@ void Algorithms::reduceVector(AccT* result, } } -#else - -namespace { -template -void launchReduction(AccT* result, const VecT* buffer, size_t size, S reducer, void* streamPtr) { - ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { - cgh.parallel_for(sycl::range<1>{size}, reducer, [=](sycl::id<1> idx, auto& redval) { - redval.combine(static_cast(buffer[idx])); - }); - }); -} -} // namespace - -namespace device { -template -void Algorithms::reduceVector(AccT* result, - const VecT* buffer, - bool overrideResult, - size_t size, - ReductionType type, - void* streamPtr) { - auto properties = [&]() -> sycl::property_list { - if (overrideResult) { - return sycl::property_list{sycl::property::reduction::initialize_to_identity()}; - } else { - return sycl::property_list{}; - } - }(); - switch (type) { - case ReductionType::Add: { - return launchReduction( - result, buffer, size, sycl::reduction(result, sycl::plus(), properties), streamPtr); - } - case ReductionType::Max: { - return launchReduction(result, - buffer, - size, - sycl::reduction(result, sycl::maximum(), properties), - streamPtr); - } - case ReductionType::Min: { - return launchReduction(result, - buffer, - size, - sycl::reduction(result, sycl::minimum(), properties), - streamPtr); - } - } -} - -#endif - template void Algorithms::reduceVector(int* result, const int* buffer, bool overrideResult, diff --git a/cuda.cmake b/cuda.cmake index 3c5e0d6..d4eabb7 100644 --- a/cuda.cmake +++ b/cuda.cmake @@ -30,16 +30,6 @@ set_source_files_properties(device.cpp string(REPLACE "sm_" "" CUDA_DEVICE_ARCH "${DEVICE_ARCH}") set_target_properties(device PROPERTIES CUDA_ARCHITECTURES "${CUDA_DEVICE_ARCH}") -target_compile_features(device PRIVATE cxx_std_17) - -target_compile_definitions(device PRIVATE $<$: - -DDEVICE_${BACKEND_UPPER_CASE}_LANG; - >) - -if (USE_GRAPH_CAPTURING) - target_compile_definitions(device PRIVATE DEVICE_USE_GRAPH_CAPTURING) -endif() - target_compile_options(device PRIVATE $<$: --expt-relaxed-constexpr; >) diff --git a/hip.cmake b/hip.cmake index f8b3d88..2171f81 100644 --- a/hip.cmake +++ b/hip.cmake @@ -45,8 +45,7 @@ if (IS_NVCC_PLATFORM) else() set(DEVICE_HIPCC -std=c++17; -O3; - --offload-arch=${DEVICE_ARCH}; - -DDEVICE_${BACKEND_UPPER_CASE}_LANG) + --offload-arch=${DEVICE_ARCH}) endif() if (DEVICE_KERNEL_INFOPRINT) diff --git a/interfaces/common/Common.h b/interfaces/common/Common.h index 6fdc465..a789209 100644 --- a/interfaces/common/Common.h +++ b/interfaces/common/Common.h @@ -8,6 +8,7 @@ #include "utils/env.h" #include "utils/logger.h" +#include #include #include #include @@ -31,25 +32,31 @@ using StatusT = std::array; template void isFlagSet(const StatusT& status) { assert(status[ID]); -}; +} template U align(T number, U alignment) { size_t alignmentFactor = (number + alignment - 1) / alignment; return alignmentFactor * alignment; } -} // namespace device -constexpr auto mapPercentage(int minval, int maxval, double value) { - if (std::isnan(value)) { - return 0; +/** + * Maps a stream priority as the device API states it - 0 the lowest, 1 the highest, NaN the + * runtime default - onto the integer range the runtime reports. + * + * CUDA and HIP report that range as a pair in which the numerically *smaller* value is the + * *higher* priority, so the mapping runs downwards from leastPriority to greatestPriority. Both + * ends coincide on devices that do not support priorities, which then yields that single value. + */ +inline int mapStreamPriority(int leastPriority, int greatestPriority, double priority) { + if (std::isnan(priority)) { + return leastPriority; } - const auto convminval = static_cast(minval); - const auto convmaxval = static_cast(maxval); - - const auto transformed = value * (convmaxval - convminval + 1) + convminval; - return std::max(std::min(static_cast(std::floor(transformed)), maxval), minval); + const auto fraction = std::min(std::max(priority, 0.0), 1.0); + const auto span = static_cast(leastPriority) - static_cast(greatestPriority); + return static_cast(std::lround(static_cast(leastPriority) - fraction * span)); } +} // namespace device #endif // SEISSOLDEVICE_INTERFACES_COMMON_COMMON_H_ diff --git a/interfaces/cuda/Control.cu b/interfaces/cuda/Control.cu index 0fe1815..408a228 100644 --- a/interfaces/cuda/Control.cu +++ b/interfaces/cuda/Control.cu @@ -5,6 +5,7 @@ #include "utils/env.h" #include "utils/logger.h" +#include #include #include #include @@ -27,7 +28,11 @@ namespace { #ifdef DEVICE_CONTEXT_GLOBAL int currentDeviceId = 0; #else -thread_local int currentDeviceId = 0; +// The runtime keeps the selected device per thread, so a thread that has not selected one works +// on device 0 - which is the wrong card whenever the process picked another one. The device the +// process selected is kept here and picked up on the first request from a thread that has none. +std::atomic selectedDeviceId{0}; +thread_local int currentDeviceId = -1; #endif } // namespace @@ -52,6 +57,9 @@ ConcreteAPI::ConcreteAPI() { void ConcreteAPI::setDevice(int deviceId) { currentDeviceId = deviceId; +#ifndef DEVICE_CONTEXT_GLOBAL + selectedDeviceId.store(deviceId, std::memory_order_relaxed); +#endif APIWRAP(cudaSetDevice(deviceId)); @@ -75,7 +83,7 @@ void ConcreteAPI::initialize() { usmDefault = properties[getDeviceId()].directManagedMemAccessFromHost != 0; - APIWRAP(cudaDeviceGetStreamPriorityRange(&priorityMin, &priorityMax)); + APIWRAP(cudaDeviceGetStreamPriorityRange(&priorityLeast, &priorityGreatest)); int canCompressProto = 0; DRVWRAP(cuDeviceGetAttribute( @@ -87,19 +95,24 @@ void ConcreteAPI::initialize() { } void ConcreteAPI::finalize() { + const std::lock_guard lock(apiMutex); if (status[StatusID::InterfaceInitialized]) { CHECK_ERR; APIWRAP(cudaStreamDestroy(defaultStream)); + defaultStream = nullptr; + if (!genericStreams.empty()) { logInfo() << "DEVICE::WARNING:" << genericStreams.size() << "device generic stream(s) were not deleted."; for (auto stream : genericStreams) { APIWRAP(cudaStreamDestroy(stream)); } + genericStreams.clear(); } status[StatusID::InterfaceInitialized] = false; } + m_isFinalized = true; } int ConcreteAPI::getNumDevices() { return properties.size(); } @@ -108,6 +121,12 @@ int ConcreteAPI::getDeviceId() { if (!status[StatusID::DeviceSelected]) { logError() << "Device has not been selected. Please, select device before requesting device Id"; } +#ifndef DEVICE_CONTEXT_GLOBAL + if (currentDeviceId < 0) { + currentDeviceId = selectedDeviceId.load(std::memory_order_relaxed); + APIWRAP(cudaSetDevice(currentDeviceId)); + } +#endif return currentDeviceId; } diff --git a/interfaces/cuda/Copy.cu b/interfaces/cuda/Copy.cu index 7776096..72adbaa 100644 --- a/interfaces/cuda/Copy.cu +++ b/interfaces/cuda/Copy.cu @@ -65,6 +65,13 @@ void ConcreteAPI::prefetchUnifiedMemTo(Destination type, size_t count, void* streamPtr) { isFlagSet(status); + + // Prefetching managed memory needs concurrent managed access, in either direction. A prefetch + // is a hint, so where the device cannot serve it, skipping is the whole handling. + if (!allowedConcurrentManagedAccess) { + return; + } + cudaStream_t stream = (streamPtr == nullptr) ? nullptr : (static_cast(streamPtr)); cudaMemLocation location{}; @@ -73,7 +80,7 @@ void ConcreteAPI::prefetchUnifiedMemTo(Destination type, #if CUDART_VERSION >= 13000 location.type = cudaMemLocationTypeHost; #endif - } else if (allowedConcurrentManagedAccess) { + } else { location.id = getDeviceId(); #if CUDART_VERSION >= 13000 location.type = cudaMemLocationTypeDevice; diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index 6233474..451f342 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -82,9 +82,18 @@ class ConcreteAPI : public AbstractAPI { void syncDefaultStreamWithHost() override; bool isCapableOfGraphCapturing() override; - DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; - void streamEndCapture(DeviceGraphHandle handle) override; - void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override; + DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) override; + void streamEndCapture(const DeviceGraphHandle& handle) override; + void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; + + bool isCapableOfGraphNodes() override; + DeviceGraphHandle graphCreate() override; + void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) override; + DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) override; + void graphInstantiate(const DeviceGraphHandle& graphHandle) override; void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; @@ -114,6 +123,9 @@ class ConcreteAPI : public AbstractAPI { void setupPrinting(int rank) override; private: + // Drops the allocation from the bookkeeping and returns its size. + size_t forgetAllocation(void* devPtr); + device::StatusT status{false}; std::vector properties; @@ -127,20 +139,13 @@ class ConcreteAPI : public AbstractAPI { std::unordered_set genericStreams{}; - struct GraphDetails { - cudaGraph_t graph; - cudaGraphExec_t instance; - std::vector streamPtrs; - bool ready{false}; - }; - std::vector graphs; - Statistics statistics{}; - std::unordered_map memToSizeMap{{nullptr, 0}}; + std::unordered_map memToSizeMap; - int priorityMin, priorityMax; + int priorityLeast{0}; + int priorityGreatest{0}; - std::unordered_map allocationProperties; + std::unordered_map allocationProperties; }; } // namespace device diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index 263e3a7..e1e6c1e 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -10,24 +10,117 @@ #include #include #include +#include +#include #include +#include +#include +#include + +// Explicit graph nodes rest on cudaStreamBeginCaptureToGraph, which the runtime gained in CUDA +// 12.3. Capturing whole streams works without it, so the two get their own macro. +#if defined(DEVICE_USE_GRAPH_CAPTURING) && (CUDART_VERSION >= 12030) +#define DEVICE_USE_GRAPH_NODES +#endif using namespace device; -/* This is a wrapped graph capturing CUDA mechanism. - * Call the following in order to capture a computational graph - * streamBeginCapture(); // 1 +/* Two ways of building a compute graph are offered. * - * // your GPU code here // 2 + * Whole-stream capture, for code that only wants to replay a fixed sequence: + * auto graph = streamBeginCapture(streams); // 1 + * // your GPU code here // 2 + * streamEndCapture(graph); // 3 + * launchGraph(graph, stream); // 4 * - * streamEndCapture(); // 3 - * auto graph = getGraphInstance(); // 4 - * - * Once you have a coompute-graph recorded you can invoke it as follows: - * launchGraph(graph) // 1 - * syncGraph(graph) // 2 + * Explicit node construction, for code that knows its own dependency structure: + * auto graph = graphCreate(); // 1 + * auto a = graphAddNode(graph, {}, stream, recordA); // 2 + * auto b = graphAddNode(graph, {a}, stream, recordB); // 3 + * graphInstantiate(graph); // 4 + * launchGraph(graph, stream); // 5 * */ +namespace { +std::mutex hostFunctionMutex; +std::unordered_map>>> + capturedHostFunctions; +} // namespace + +namespace device::internals { +CaptureState captureState(cudaStream_t stream) { + CaptureState state{}; + unsigned long long captureId{}; + const cudaGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + + // The unversioned name resolves to different signatures depending on the toolkit: up to CUDA + // 12.x it is the six-argument form, from CUDA 13 on it is the one that also reports edge data. + // cudaStreamGetCaptureInfo_v2 is not an option, as CUDA 13 no longer declares it. +#if CUDART_VERSION >= 13000 + const cudaGraphEdgeData* edgeData{nullptr}; + APIWRAP(cudaStreamGetCaptureInfo( + stream, &state.status, &captureId, &state.graph, &frontier, &edgeData, &frontierSize)); +#else + APIWRAP(cudaStreamGetCaptureInfo( + stream, &state.status, &captureId, &state.graph, &frontier, &frontierSize)); +#endif + + if (frontier != nullptr) { + state.frontier.assign(frontier, frontier + frontierSize); + } + return state; +} + +std::function* adoptHostFunction(cudaGraph_t graph, const std::function& function) { + if (graph == nullptr) { + return nullptr; + } + + const std::lock_guard lock(hostFunctionMutex); + auto& functions = capturedHostFunctions[graph]; + functions.emplace_back(std::make_unique>(function)); + return functions.back().get(); +} + +void forgetHostFunctions(cudaGraph_t graph) { + const std::lock_guard lock(hostFunctionMutex); + capturedHostFunctions.erase(graph); +} +} // namespace device::internals + +namespace device { +struct DeviceGraph { + cudaGraph_t graph{nullptr}; + cudaGraphExec_t instance{nullptr}; + + // one entry per graphAddNode call; an entry may hold zero, one or several native nodes + std::vector> nodes; + + // only used by the whole-stream capture path + std::vector streamPtrs; + + bool ready{false}; + + DeviceGraph() = default; + DeviceGraph(const DeviceGraph&) = delete; + DeviceGraph& operator=(const DeviceGraph&) = delete; + + ~DeviceGraph() { + internals::forgetHostFunctions(graph); + + // deliberately unchecked: the graph may outlive the device context during teardown, and a + // failure here has nothing left to report to + if (instance != nullptr) { + cudaGraphExecDestroy(instance); + } + if (graph != nullptr) { + cudaGraphDestroy(graph); + } + } +}; +} // namespace device + bool ConcreteAPI::isCapableOfGraphCapturing() { #ifdef DEVICE_USE_GRAPH_CAPTURING return true; @@ -36,55 +129,122 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { #endif } -DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { - auto handle = DeviceGraphHandle(); +bool ConcreteAPI::isCapableOfGraphNodes() { +#ifdef DEVICE_USE_GRAPH_NODES + return true; +#else + return false; +#endif +} + +DeviceGraphHandle ConcreteAPI::streamBeginCapture(const std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING - { - std::lock_guard guard(apiMutex); - graphs.push_back(GraphDetails{}); - handle = DeviceGraphHandle(graphs.size() - 1); - - GraphDetails& graphInstance = graphs[handle.getGraphId()]; - graphInstance.ready = false; - graphInstance.streamPtrs = streamPtrs; + if (streamPtrs.empty()) { + logError() << "Graph capturing records streams, so it needs at least one."; + return DeviceGraphHandle(); } + auto graphInstance = std::make_shared(); + graphInstance->streamPtrs = streamPtrs; + APIWRAP(cudaStreamBeginCapture(static_cast(streamPtrs[0]), cudaStreamCaptureModeThreadLocal)); + + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); #endif - return handle; } -void ConcreteAPI::streamEndCapture(DeviceGraphHandle handle) { +void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[handle.getGraphId()]; - } - APIWRAP(cudaStreamEndCapture(static_cast(graphInstance.streamPtrs[0]), - &(graphInstance.graph))); + auto* graphInstance = handle.get(); + assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); + + APIWRAP(cudaStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), + &(graphInstance->graph))); APIWRAP( - cudaGraphInstantiate(&(graphInstance.instance), graphInstance.graph, nullptr, nullptr, 0)); + cudaGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); - graphInstance.ready = true; + graphInstance->ready = true; +#endif +} - { - std::lock_guard guard(apiMutex); - graphs[handle.getGraphId()] = graphInstance; +DeviceGraphHandle ConcreteAPI::graphCreate() { +#ifdef DEVICE_USE_GRAPH_NODES + auto graphInstance = std::make_shared(); + APIWRAP(cudaGraphCreate(&(graphInstance->graph), 0)); + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); +#endif +} + +void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_NODES + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); + assert(!graphInstance->ready && "no nodes can be added to an instantiated graph"); + + std::vector nativeDependencies; + for (const auto& dependency : dependencies) { + assert(dependency.isInitialized() && "an uninitialized node cannot be depended upon"); + const auto& nodes = graphInstance->nodes.at(dependency.getNodeId()); + nativeDependencies.insert(nativeDependencies.end(), nodes.begin(), nodes.end()); } + + APIWRAP(cudaStreamBeginCaptureToGraph(static_cast(streamPtr), + graphInstance->graph, + nativeDependencies.data(), + nullptr, + nativeDependencies.size(), + cudaStreamCaptureModeThreadLocal)); #endif } -void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) { +DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_NODES + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a node must be opened before it can be closed"); + + auto stream = static_cast(streamPtr); + auto produced = internals::captureState(stream).frontier; + + cudaGraph_t endedGraph{nullptr}; + APIWRAP(cudaStreamEndCapture(stream, &endedGraph)); + assert(endedGraph == graphInstance->graph && "capturing into a graph hands that same graph back"); + + graphInstance->nodes.emplace_back(std::move(produced)); + return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); +#else + return DeviceGraphNodeHandle(); +#endif +} + +void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { +#ifdef DEVICE_USE_GRAPH_NODES + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); + + APIWRAP( + cudaGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); + + graphInstance->ready = true; +#endif +} + +void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING - assert(graphHandle.isInitialized() && "a graph must be captured before launching"); - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[graphHandle.getGraphId()]; - } - APIWRAP(cudaGraphLaunch(graphInstance.instance, reinterpret_cast(streamPtr))); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && graphInstance->ready && + "a graph must be captured before launching"); + + APIWRAP(cudaGraphLaunch(graphInstance->instance, static_cast(streamPtr))); #endif } diff --git a/interfaces/cuda/Internals.cu b/interfaces/cuda/Internals.cu index 419b27f..44c54e3 100644 --- a/interfaces/cuda/Internals.cu +++ b/interfaces/cuda/Internals.cu @@ -4,26 +4,27 @@ #include "utils/logger.h" +#include #include +#include #include #include -#include namespace device::internals { -thread_local std::string prevFile{}; +thread_local const char* prevFile{nullptr}; thread_local int prevLine{-1}; cudaError_t checkResult(cudaError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except) { - if (error != cudaSuccess && except.find(error) == except.end()) { + std::initializer_list except) { + if (error != cudaSuccess && std::find(except.begin(), except.end(), error) == except.end()) { std::stringstream stream; stream << '\n' << file << ", line " << line << ": " << cudaGetErrorString(error) << " (" << error << ")\n"; - if (prevLine >= 0) { + if (prevFile != nullptr) { stream << "Previous CUDA API/Driver call:" << std::endl << prevFile << ", line " << prevLine << std::endl; } @@ -35,10 +36,10 @@ cudaError_t checkResult(cudaError_t error, } CUresult checkResultDriver(CUresult error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except) { - if (error != CUDA_SUCCESS && except.find(error) == except.end()) { + std::initializer_list except) { + if (error != CUDA_SUCCESS && std::find(except.begin(), except.end(), error) == except.end()) { const char* errstr = nullptr; const auto errstrRes = cuGetErrorString(error, &errstr); @@ -50,7 +51,7 @@ CUresult checkResultDriver(CUresult error, stream << "[ERROR WHILE RETRIEVING ERROR STRING]"; } stream << " (" << error << ")\n"; - if (prevLine >= 0) { + if (prevFile != nullptr) { stream << "Previous CUDA API/Driver call:" << std::endl << prevFile << ", line " << prevLine << std::endl; } diff --git a/interfaces/cuda/Internals.h b/interfaces/cuda/Internals.h index 497c739..7918944 100644 --- a/interfaces/cuda/Internals.h +++ b/interfaces/cuda/Internals.h @@ -6,8 +6,9 @@ #define SEISSOLDEVICE_INTERFACES_CUDA_INTERNALS_H_ #include -#include -#include +#include +#include +#include #include #define APIWRAP(call) (void)::device::internals::checkResult(call, __FILE__, __LINE__, {}) @@ -20,16 +21,39 @@ namespace device::internals { using DeviceStreamT = cudaStream_t; +/** + * What a stream is currently recording into: the capture status, the graph the operations end up + * in, and the nodes a subsequently recorded operation would depend on. + */ +struct CaptureState { + cudaStreamCaptureStatus status{}; + cudaGraph_t graph{nullptr}; + std::vector frontier; +}; + +CaptureState captureState(cudaStream_t stream); + +/** + * Hands a host function to the graph that is being recorded, which keeps it alive for as long as + * the graph can be replayed, and returns the copy to pass to the runtime. Returns nullptr if the + * graph is not one of ours. + */ +std::function* adoptHostFunction(cudaGraph_t graph, const std::function& function); +void forgetHostFunctions(cudaGraph_t graph); + constexpr static int DefaultBlockDim = 1024; +// Every wrapped call goes through here, so the parameters stay free of anything that allocates: +// the file name is the string literal __FILE__ expands to, and the accepted errors are read from +// the caller's temporary array. cudaError_t checkResult(cudaError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except); + std::initializer_list except); CUresult checkResultDriver(CUresult error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except); + std::initializer_list except); inline dim3 computeGrid1D(const dim3& block, const size_t size) { int numBlocks = (size + block.x - 1) / block.x; diff --git a/interfaces/cuda/Memory.cu b/interfaces/cuda/Memory.cu index df55a59..6587525 100644 --- a/interfaces/cuda/Memory.cu +++ b/interfaces/cuda/Memory.cu @@ -12,6 +12,7 @@ #include #include #include +#include #include using namespace device; @@ -66,6 +67,7 @@ void driverFree(void* ptr, std::size_t size, const CUmemAllocationProp& prop) { void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr = nullptr; if (compress && canCompress) { CUmemAllocationProp prop = {}; @@ -76,7 +78,7 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC; devPtr = driverAllocate(size, prop); - allocationProperties[devPtr] = reinterpret_cast(new CUmemAllocationProp(prop)); + allocationProperties[devPtr] = prop; } else { APIWRAP(cudaMalloc(&devPtr, size)); } @@ -87,31 +89,39 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr = nullptr; APIWRAP(cudaMallocManaged(&devPtr, size, cudaMemAttachGlobal)); - cudaMemLocation location{}; - if (hint == Destination::Host) { - location.id = cudaCpuDeviceId; + // Naming the device as the preferred location needs concurrent managed access. Where that is + // missing there is no location to name instead, so the allocation keeps the driver default - + // passing a zeroed location would advise for device 0 up to CUDA 12 and be rejected from CUDA + // 13 on. + const bool hasPreferredLocation = (hint == Destination::Host) || allowedConcurrentManagedAccess; + if (hasPreferredLocation) { + cudaMemLocation location{}; + if (hint == Destination::Host) { + location.id = cudaCpuDeviceId; #if CUDART_VERSION >= 13000 - location.type = cudaMemLocationTypeHost; + location.type = cudaMemLocationTypeHost; #endif - } else if (allowedConcurrentManagedAccess) { - location.id = getDeviceId(); + } else { + location.id = getDeviceId(); #if CUDART_VERSION >= 13000 - location.type = cudaMemLocationTypeDevice; + location.type = cudaMemLocationTypeDevice; #endif - } + } - APIWRAP(cudaMemAdvise(devPtr, - size, - cudaMemAdviseSetPreferredLocation, + APIWRAP(cudaMemAdvise(devPtr, + size, + cudaMemAdviseSetPreferredLocation, #if CUDART_VERSION >= 13000 - location + location #else - location.id + location.id #endif - )); + )); + } statistics.allocatedMemBytes += size; statistics.allocatedUnifiedMemBytes += size; @@ -121,6 +131,7 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr = nullptr; const auto flag = hint == Destination::Host ? cudaHostAllocDefault : cudaHostAllocMapped; APIWRAP(cudaHostAlloc(&devPtr, size, flag)); @@ -129,15 +140,35 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) return devPtr; } +size_t ConcreteAPI::forgetAllocation(void* devPtr) { + const auto entry = memToSizeMap.find(devPtr); + if (entry == memToSizeMap.end()) { + assert(false && "DEVICE: an attempt to delete mem. which has not been allocated. unknown " + "pointer"); + return 0; + } + + const auto size = entry->second; + memToSizeMap.erase(entry); + statistics.deallocatedMemBytes += size; + return size; +} + void ConcreteAPI::freeGlobMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; - if (allocationProperties.find(devPtr) != allocationProperties.end()) { - driverFree(devPtr, - memToSizeMap.at(devPtr), - *reinterpret_cast(allocationProperties.at(devPtr))); + const std::lock_guard lock(apiMutex); + if (devPtr == nullptr) { + return; + } + + const auto size = forgetAllocation(devPtr); + + const auto properties = allocationProperties.find(devPtr); + if (properties != allocationProperties.end()) { + driverFree(devPtr, size, properties->second); + // the entry has to go with the allocation: the runtime is free to hand the same address out + // again, and a leftover entry would send that one down the driver path as well + allocationProperties.erase(properties); } else { APIWRAP(cudaFree(devPtr)); } @@ -145,17 +176,24 @@ void ConcreteAPI::freeGlobMem(void* devPtr) { void ConcreteAPI::freeUnifiedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + const std::lock_guard lock(apiMutex); + if (devPtr == nullptr) { + return; + } + + const auto size = forgetAllocation(devPtr); + statistics.allocatedUnifiedMemBytes -= size; APIWRAP(cudaFree(devPtr)); } void ConcreteAPI::freePinnedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + const std::lock_guard lock(apiMutex); + if (devPtr == nullptr) { + return; + } + + forgetAllocation(devPtr); APIWRAP(cudaFreeHost(devPtr)); } @@ -176,6 +214,7 @@ void ConcreteAPI::freeMemAsync(void* devPtr, void* streamPtr) { std::string ConcreteAPI::getMemLeaksReport() { isFlagSet(status); + const std::lock_guard lock(apiMutex); std::ostringstream report{}; report << "Memory Leaks, bytes: " << (statistics.allocatedMemBytes - statistics.deallocatedMemBytes) << '\n'; @@ -186,11 +225,13 @@ size_t ConcreteAPI::getMaxAvailableMem() { return properties[getDeviceId()].tota size_t ConcreteAPI::getCurrentlyOccupiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedMemBytes; } size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedUnifiedMemBytes; } diff --git a/interfaces/cuda/Streams.cu b/interfaces/cuda/Streams.cu index b70156c..6e6d69f 100644 --- a/interfaces/cuda/Streams.cu +++ b/interfaces/cuda/Streams.cu @@ -9,6 +9,7 @@ #include #include #include +#include #include using namespace device; @@ -28,8 +29,9 @@ void ConcreteAPI::syncDefaultStreamWithHost() { void* ConcreteAPI::createStream(double priority) { isFlagSet(status); + const std::lock_guard lock(apiMutex); cudaStream_t stream; - const auto truePriority = mapPercentage(priorityMin, priorityMax, priority); + const auto truePriority = mapStreamPriority(priorityLeast, priorityGreatest, priority); APIWRAP(cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, truePriority)); genericStreams.insert(stream); return reinterpret_cast(stream); @@ -37,11 +39,20 @@ void* ConcreteAPI::createStream(double priority) { void ConcreteAPI::destroyGenericStream(void* streamPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); cudaStream_t stream = static_cast(streamPtr); + + // The stream has to leave the set before it is destroyed, and a stream that is not in it is not + // this backend's to destroy - the default stream, for one, would take the whole interface with + // it. auto it = genericStreams.find(stream); - if (it != genericStreams.end()) { - genericStreams.erase(it); + if (it == genericStreams.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or was not created here; not destroying it."; + return; } + + genericStreams.erase(it); APIWRAP(cudaStreamDestroy(stream)); } @@ -70,13 +81,15 @@ void ConcreteAPI::syncStreamWithEvent(void* streamPtr, void* eventPtr) { } namespace { +// Called once, so the copy goes away with the call. void streamCallbackEpheremal(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); delete function; } -void streamCallbackPermanent(void* data) { +// Called on every replay of the graph it was recorded into, which owns the copy. +void streamCallbackRecorded(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); } @@ -85,17 +98,23 @@ void streamCallbackPermanent(void* data) { void ConcreteAPI::streamHostFunction(void* streamPtr, const std::function& function) { cudaStream_t stream = static_cast(streamPtr); - cudaStreamCaptureStatus status{}; - APIWRAP(cudaStreamIsCapturing(stream, &status)); + const auto capture = internals::captureState(stream); + if (capture.status == cudaStreamCaptureStatusInvalidated) { + return; + } - if (status != cudaStreamCaptureStatusInvalidated) { - auto* functionData = new std::function(function); - if (status == cudaStreamCaptureStatusActive) { - APIWRAP(cudaLaunchHostFunc(stream, &streamCallbackPermanent, functionData)); - } else { - APIWRAP(cudaLaunchHostFunc(stream, &streamCallbackEpheremal, functionData)); + if (capture.status == cudaStreamCaptureStatusActive) { + auto* recorded = internals::adoptHostFunction(capture.graph, function); + if (recorded == nullptr) { + logError() << "A host function was recorded into a graph this backend does not know."; + return; } + APIWRAP(cudaLaunchHostFunc(stream, &streamCallbackRecorded, recorded)); + return; } + + APIWRAP( + cudaLaunchHostFunc(stream, &streamCallbackEpheremal, new std::function(function))); } namespace { diff --git a/interfaces/hip/Control.cpp b/interfaces/hip/Control.cpp index 839c557..bffd63e 100644 --- a/interfaces/hip/Control.cpp +++ b/interfaces/hip/Control.cpp @@ -6,6 +6,7 @@ #include "utils/env.h" #include "utils/logger.h" +#include #include #include #include @@ -49,6 +50,9 @@ ConcreteAPI::ConcreteAPI() { void ConcreteAPI::setDevice(int deviceId) { currentDeviceId = deviceId; +#ifndef DEVICE_CONTEXT_GLOBAL + selectedDeviceId.store(deviceId, std::memory_order_relaxed); +#endif APIWRAP(hipSetDevice(deviceId)); @@ -84,27 +88,31 @@ void ConcreteAPI::initialize() { properties[getDeviceId()].pageableMemoryAccessUsesHostPageTables != 0; } - APIWRAP(hipDeviceGetStreamPriorityRange(&priorityMin, &priorityMax)); + APIWRAP(hipDeviceGetStreamPriorityRange(&priorityLeast, &priorityGreatest)); } else { logWarning() << "Device Interface has already been initialized"; } } void ConcreteAPI::finalize() { + const std::lock_guard lock(apiMutex); if (status[StatusID::InterfaceInitialized]) { - CHECK_ERR; APIWRAP(hipStreamDestroy(defaultStream)); + defaultStream = nullptr; + if (!genericStreams.empty()) { logInfo() << "DEVICE::WARNING:" << genericStreams.size() << "device generic stream(s) were not deleted."; for (auto stream : genericStreams) { APIWRAP(hipStreamDestroy(stream)); } + genericStreams.clear(); } status[StatusID::InterfaceInitialized] = false; } + m_isFinalized = true; } int ConcreteAPI::getNumDevices() { return properties.size(); } @@ -113,6 +121,12 @@ int ConcreteAPI::getDeviceId() { if (!status[StatusID::DeviceSelected]) { logError() << "Device has not been selected. Please, select device before requesting device Id"; } +#ifndef DEVICE_CONTEXT_GLOBAL + if (currentDeviceId < 0) { + currentDeviceId = selectedDeviceId.load(std::memory_order_relaxed); + APIWRAP(hipSetDevice(currentDeviceId)); + } +#endif return currentDeviceId; } diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index c7a5cf4..3d2ec61 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -8,23 +8,111 @@ #include "utils/logger.h" #include +#include +#include +#include +#include +#include +#include +#include +#include + +// Explicit graph nodes rest on hipStreamBeginCaptureToGraph, which HIP gained in ROCm 6.3. +// Capturing whole streams works without it, so the two get their own macro. +#if defined(DEVICE_USE_GRAPH_CAPTURING) && \ + (HIP_VERSION_MAJOR > 6 || (HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR >= 3)) +#define DEVICE_USE_GRAPH_NODES +#endif using namespace device; -/* This is a wrapped graph capturing CUDA mechanism. - * Call the following in order to capture a computational graph - * streamBeginCapture(); // 1 +/* Two ways of building a compute graph are offered. * - * // your GPU code here // 2 + * Whole-stream capture, for code that only wants to replay a fixed sequence: + * auto graph = streamBeginCapture(streams); // 1 + * // your GPU code here // 2 + * streamEndCapture(graph); // 3 + * launchGraph(graph, stream); // 4 * - * streamEndCapture(); // 3 - * auto graph = getGraphInstance(); // 4 - * - * Once you have a coompute-graph recorded you can invoke it as follows: - * launchGraph(graph) // 1 - * syncGraph(graph) // 2 + * Explicit node construction, for code that knows its own dependency structure: + * auto graph = graphCreate(); // 1 + * auto a = graphAddNode(graph, {}, stream, recordA); // 2 + * auto b = graphAddNode(graph, {a}, stream, recordB); // 3 + * graphInstantiate(graph); // 4 + * launchGraph(graph, stream); // 5 * */ +namespace { +std::mutex hostFunctionMutex; +std::unordered_map>>> + capturedHostFunctions; +} // namespace + +namespace device::internals { +CaptureState captureState(hipStream_t stream) { + CaptureState state{}; + unsigned long long captureId{}; + const hipGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + + APIWRAP(hipStreamGetCaptureInfo_v2( + stream, &state.status, &captureId, &state.graph, &frontier, &frontierSize)); + + if (frontier != nullptr) { + state.frontier.assign(frontier, frontier + frontierSize); + } + return state; +} + +std::function* adoptHostFunction(hipGraph_t graph, const std::function& function) { + if (graph == nullptr) { + return nullptr; + } + + const std::lock_guard lock(hostFunctionMutex); + auto& functions = capturedHostFunctions[graph]; + functions.emplace_back(std::make_unique>(function)); + return functions.back().get(); +} + +void forgetHostFunctions(hipGraph_t graph) { + const std::lock_guard lock(hostFunctionMutex); + capturedHostFunctions.erase(graph); +} +} // namespace device::internals + +namespace device { +struct DeviceGraph { + hipGraph_t graph{nullptr}; + hipGraphExec_t instance{nullptr}; + + // one entry per graphAddNode call; an entry may hold zero, one or several native nodes + std::vector> nodes; + + // only used by the whole-stream capture path + std::vector streamPtrs; + + bool ready{false}; + + DeviceGraph() = default; + DeviceGraph(const DeviceGraph&) = delete; + DeviceGraph& operator=(const DeviceGraph&) = delete; + + ~DeviceGraph() { + internals::forgetHostFunctions(graph); + + // deliberately unchecked: the graph may outlive the device context during teardown, and a + // failure here has nothing left to report to + if (instance != nullptr) { + hipGraphExecDestroy(instance); + } + if (graph != nullptr) { + hipGraphDestroy(graph); + } + } +}; +} // namespace device + bool ConcreteAPI::isCapableOfGraphCapturing() { #ifdef DEVICE_USE_GRAPH_CAPTURING return true; @@ -33,54 +121,123 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { #endif } -DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { - auto handle = DeviceGraphHandle(); +bool ConcreteAPI::isCapableOfGraphNodes() { +#ifdef DEVICE_USE_GRAPH_NODES + return true; +#else + return false; +#endif +} + +DeviceGraphHandle ConcreteAPI::streamBeginCapture(const std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING - { - std::lock_guard guard(apiMutex); - graphs.push_back(GraphDetails{}); - handle = DeviceGraphHandle(graphs.size() - 1); - - GraphDetails& graphInstance = graphs[handle.getGraphId()]; - graphInstance.ready = false; - graphInstance.streamPtrs = streamPtrs; + if (streamPtrs.empty()) { + logError() << "Graph capturing records streams, so it needs at least one."; + return DeviceGraphHandle(); } + auto graphInstance = std::make_shared(); + graphInstance->streamPtrs = streamPtrs; + APIWRAP(hipStreamBeginCapture(static_cast(streamPtrs[0]), hipStreamCaptureModeThreadLocal)); + + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); #endif - return handle; } -void ConcreteAPI::streamEndCapture(DeviceGraphHandle handle) { +void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[handle.getGraphId()]; - } - APIWRAP(hipStreamEndCapture(static_cast(graphInstance.streamPtrs[0]), - &(graphInstance.graph))); + auto* graphInstance = handle.get(); + assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); + + APIWRAP(hipStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), + &(graphInstance->graph))); + + APIWRAP( + hipGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); - APIWRAP(hipGraphInstantiate(&(graphInstance.instance), graphInstance.graph, nullptr, nullptr, 0)); + graphInstance->ready = true; +#endif +} - graphInstance.ready = true; +DeviceGraphHandle ConcreteAPI::graphCreate() { +#ifdef DEVICE_USE_GRAPH_NODES + auto graphInstance = std::make_shared(); + APIWRAP(hipGraphCreate(&(graphInstance->graph), 0)); + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); +#endif +} + +void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_NODES + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); + assert(!graphInstance->ready && "no nodes can be added to an instantiated graph"); - { - std::lock_guard guard(apiMutex); - graphs[handle.getGraphId()] = graphInstance; + std::vector nativeDependencies; + for (const auto& dependency : dependencies) { + assert(dependency.isInitialized() && "an uninitialized node cannot be depended upon"); + const auto& nodes = graphInstance->nodes.at(dependency.getNodeId()); + nativeDependencies.insert(nativeDependencies.end(), nodes.begin(), nodes.end()); } + + // the edge-data argument is not supported by HIP and has to stay a nullptr + APIWRAP(hipStreamBeginCaptureToGraph(static_cast(streamPtr), + graphInstance->graph, + nativeDependencies.data(), + nullptr, + nativeDependencies.size(), + hipStreamCaptureModeThreadLocal)); #endif } -void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) { +DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_NODES + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a node must be opened before it can be closed"); + + auto stream = static_cast(streamPtr); + auto produced = internals::captureState(stream).frontier; + + hipGraph_t endedGraph{nullptr}; + APIWRAP(hipStreamEndCapture(stream, &endedGraph)); + assert(endedGraph == graphInstance->graph && "capturing into a graph hands that same graph back"); + + graphInstance->nodes.emplace_back(std::move(produced)); + return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); +#else + return DeviceGraphNodeHandle(); +#endif +} + +void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { +#ifdef DEVICE_USE_GRAPH_NODES + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); + + APIWRAP( + hipGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); + + graphInstance->ready = true; +#endif +} + +void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING - assert(graphHandle.isInitialized() && "a graph must be captured before launching"); - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[graphHandle.getGraphId()]; - } - APIWRAP(hipGraphLaunch(graphInstance.instance, reinterpret_cast(streamPtr))); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && graphInstance->ready && + "a graph must be captured before launching"); + + APIWRAP(hipGraphLaunch(graphInstance->instance, static_cast(streamPtr))); #endif } diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index 121b55b..953b8b6 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -81,9 +81,18 @@ class ConcreteAPI : public AbstractAPI { void syncDefaultStreamWithHost() override; bool isCapableOfGraphCapturing() override; - DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; - void streamEndCapture(DeviceGraphHandle handle) override; - void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override; + DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) override; + void streamEndCapture(const DeviceGraphHandle& handle) override; + void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; + + bool isCapableOfGraphNodes() override; + DeviceGraphHandle graphCreate() override; + void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) override; + DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) override; + void graphInstantiate(const DeviceGraphHandle& graphHandle) override; void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; @@ -113,6 +122,9 @@ class ConcreteAPI : public AbstractAPI { void setupPrinting(int rank) override; private: + // Drops the allocation from the bookkeeping and returns its size. + size_t forgetAllocation(void* devPtr); + device::StatusT status{false}; std::vector properties; @@ -123,18 +135,11 @@ class ConcreteAPI : public AbstractAPI { std::unordered_set genericStreams{}; - struct GraphDetails { - hipGraph_t graph; - hipGraphExec_t instance; - std::vector streamPtrs; - bool ready{false}; - }; - std::vector graphs; - Statistics statistics{}; - std::unordered_map memToSizeMap{{nullptr, 0}}; + std::unordered_map memToSizeMap; - int priorityMin, priorityMax; + int priorityLeast{0}; + int priorityGreatest{0}; }; } // namespace device diff --git a/interfaces/hip/Internals.cpp b/interfaces/hip/Internals.cpp index 7be2d80..a435e03 100644 --- a/interfaces/hip/Internals.cpp +++ b/interfaces/hip/Internals.cpp @@ -6,25 +6,26 @@ #include "utils/logger.h" +#include +#include #include -#include #include namespace device::internals { -thread_local std::string prevFile{}; +thread_local const char* prevFile{nullptr}; thread_local int prevLine{0}; hipError_t checkResult(hipError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except) { - if (error != hipSuccess && except.find(error) == except.end()) { + std::initializer_list except) { + if (error != hipSuccess && std::find(except.begin(), except.end(), error) == except.end()) { std::stringstream stream; stream << '\n' << file << ", line " << line << ": " << hipGetErrorString(error) << " (" << error << ")\n"; - if (prevLine > 0) { + if (prevFile != nullptr) { stream << "Previous HIP call:" << std::endl << prevFile << ", line " << prevLine << std::endl; } logError() << stream.str(); diff --git a/interfaces/hip/Internals.h b/interfaces/hip/Internals.h index a3ececd..7499e73 100644 --- a/interfaces/hip/Internals.h +++ b/interfaces/hip/Internals.h @@ -7,8 +7,9 @@ #include "hip/hip_runtime.h" -#include -#include +#include +#include +#include #define APIWRAP(call) (void)::device::internals::checkResult(call, __FILE__, __LINE__, {}) #define APIWRAPX(call, except) ::device::internals::checkResult(call, __FILE__, __LINE__, except) @@ -19,10 +20,33 @@ namespace device::internals { constexpr static int DefaultBlockDim = 1024; using DeviceStreamT = hipStream_t; + +/** + * What a stream is currently recording into: the capture status, the graph the operations end up + * in, and the nodes a subsequently recorded operation would depend on. + */ +struct CaptureState { + hipStreamCaptureStatus status{}; + hipGraph_t graph{nullptr}; + std::vector frontier; +}; + +CaptureState captureState(hipStream_t stream); + +/** + * Hands a host function to the graph that is being recorded, which keeps it alive for as long as + * the graph can be replayed, and returns the copy to pass to the runtime. Returns nullptr if the + * graph is not one of ours. + */ +std::function* adoptHostFunction(hipGraph_t graph, const std::function& function); +void forgetHostFunctions(hipGraph_t graph); +// Every wrapped call goes through here, so the parameters stay free of anything that allocates: +// the file name is the string literal __FILE__ expands to, and the accepted errors are read from +// the caller's temporary array. hipError_t checkResult(hipError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except); + std::initializer_list except); inline dim3 computeGrid1D(const dim3& block, const size_t size) { int numBlocks = (size + block.x - 1) / block.x; return dim3(numBlocks, 1, 1); diff --git a/interfaces/hip/Memory.cpp b/interfaces/hip/Memory.cpp index b4ef78a..b5fa2e5 100644 --- a/interfaces/hip/Memory.cpp +++ b/interfaces/hip/Memory.cpp @@ -9,12 +9,14 @@ #include #include +#include #include using namespace device; void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr; APIWRAP(hipMalloc(&devPtr, size)); statistics.allocatedMemBytes += size; @@ -24,11 +26,12 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr; APIWRAP(hipMallocManaged(&devPtr, size, hipMemAttachGlobal)); // make coarse-grained memory access behavior the default (match with allocGlobMem) - APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetCoarseGrain, 1)); + APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetCoarseGrain, getDeviceId())); if (hint == Destination::Host) { APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetPreferredLocation, hipCpuDeviceId)); @@ -44,6 +47,7 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr; const auto flag = hint == Destination::Host ? hipHostMallocDefault : hipHostMallocMapped; APIWRAP(hipHostMalloc(&devPtr, size, flag)); @@ -52,27 +56,51 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) return devPtr; } +size_t ConcreteAPI::forgetAllocation(void* devPtr) { + const auto entry = memToSizeMap.find(devPtr); + if (entry == memToSizeMap.end()) { + assert(false && "DEVICE: an attempt to delete mem. which has not been allocated. unknown " + "pointer"); + return 0; + } + + const auto size = entry->second; + memToSizeMap.erase(entry); + statistics.deallocatedMemBytes += size; + return size; +} + void ConcreteAPI::freeGlobMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + const std::lock_guard lock(apiMutex); + if (devPtr == nullptr) { + return; + } + + forgetAllocation(devPtr); APIWRAP(hipFree(devPtr)); } void ConcreteAPI::freeUnifiedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + const std::lock_guard lock(apiMutex); + if (devPtr == nullptr) { + return; + } + + const auto size = forgetAllocation(devPtr); + statistics.allocatedUnifiedMemBytes -= size; APIWRAP(hipFree(devPtr)); } void ConcreteAPI::freePinnedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + const std::lock_guard lock(apiMutex); + if (devPtr == nullptr) { + return; + } + + forgetAllocation(devPtr); APIWRAP(hipHostFree(devPtr)); } @@ -93,6 +121,7 @@ void ConcreteAPI::freeMemAsync(void* devPtr, void* streamPtr) { std::string ConcreteAPI::getMemLeaksReport() { isFlagSet(status); + const std::lock_guard lock(apiMutex); std::ostringstream report{}; report << "Memory Leaks, bytes: " << (statistics.allocatedMemBytes - statistics.deallocatedMemBytes) << '\n'; @@ -103,11 +132,13 @@ size_t ConcreteAPI::getMaxAvailableMem() { return properties[getDeviceId()].tota size_t ConcreteAPI::getCurrentlyOccupiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedMemBytes; } size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedUnifiedMemBytes; } diff --git a/interfaces/hip/Streams.cpp b/interfaces/hip/Streams.cpp index 37a63e1..5045691 100644 --- a/interfaces/hip/Streams.cpp +++ b/interfaces/hip/Streams.cpp @@ -8,6 +8,7 @@ #include #include +#include #include using namespace device; @@ -27,20 +28,30 @@ void ConcreteAPI::syncDefaultStreamWithHost() { void* ConcreteAPI::createStream(double priority) { isFlagSet(status); + const std::lock_guard lock(apiMutex); hipStream_t stream; - const auto truePriority = mapPercentage(priorityMin, priorityMax, priority); - APIWRAP(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority)); + const auto truePriority = mapStreamPriority(priorityLeast, priorityGreatest, priority); + APIWRAP(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, truePriority)); genericStreams.insert(stream); return reinterpret_cast(stream); } void ConcreteAPI::destroyGenericStream(void* streamPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); hipStream_t stream = static_cast(streamPtr); + + // The stream has to leave the set before it is destroyed, and a stream that is not in it is not + // this backend's to destroy - the default stream, for one, would take the whole interface with + // it. auto it = genericStreams.find(stream); - if (it != genericStreams.end()) { - genericStreams.erase(it); + if (it == genericStreams.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or was not created here; not destroying it."; + return; } + + genericStreams.erase(it); APIWRAP(hipStreamDestroy(stream)); } @@ -69,13 +80,14 @@ void ConcreteAPI::syncStreamWithEvent(void* streamPtr, void* eventPtr) { } namespace { +// Called once, so the copy goes away with the call. void streamCallbackEpheremal(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); delete function; } -void streamCallbackPermanent(void* data) { +void streamCallbackRecorded(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); } @@ -84,17 +96,22 @@ void streamCallbackPermanent(void* data) { void ConcreteAPI::streamHostFunction(void* streamPtr, const std::function& function) { hipStream_t stream = static_cast(streamPtr); - hipStreamCaptureStatus status{}; - APIWRAP(hipStreamIsCapturing(stream, &status)); + const auto capture = internals::captureState(stream); + if (capture.status == hipStreamCaptureStatusInvalidated) { + return; + } - if (status != hipStreamCaptureStatusInvalidated) { - auto* functionData = new std::function(function); - if (status == hipStreamCaptureStatusActive) { - APIWRAP(hipLaunchHostFunc(stream, &streamCallbackPermanent, functionData)); - } else { - APIWRAP(hipLaunchHostFunc(stream, &streamCallbackEpheremal, functionData)); + if (capture.status == hipStreamCaptureStatusActive) { + auto* recorded = internals::adoptHostFunction(capture.graph, function); + if (recorded == nullptr) { + logError() << "A host function was recorded into a graph this backend does not know."; + return; } + APIWRAP(hipLaunchHostFunc(stream, &streamCallbackRecorded, recorded)); + return; } + + APIWRAP(hipLaunchHostFunc(stream, &streamCallbackEpheremal, new std::function(function))); } namespace { diff --git a/interfaces/sycl/Control.cpp b/interfaces/sycl/Control.cpp index aa1ada0..2402a42 100644 --- a/interfaces/sycl/Control.cpp +++ b/interfaces/sycl/Control.cpp @@ -7,7 +7,9 @@ #include "SyclWrappedAPI.h" #include "utils/logger.h" +#include #include +#include #include #include @@ -24,7 +26,7 @@ using namespace device; void ConcreteAPI::initDevices() { if (this->deviceInitialized) { - throw new std::invalid_argument("Cannot initialize the devices twice!"); + throw std::invalid_argument("Cannot initialize the devices twice!"); } for (const auto& platform : sycl::platform::get_platforms()) { @@ -42,17 +44,19 @@ void ConcreteAPI::initDevices() { } } - DeviceContext* context = new DeviceContext{device, 1}; + DeviceContext* context = new DeviceContext{device}; this->availableDevices.push_back(context); } } - sort(this->availableDevices.begin(), - this->availableDevices.end(), - [&](DeviceContext* c1, DeviceContext* c2) { - return compare(c1->queueBuffer.getDefaultQueue().get_device(), - c2->queueBuffer.getDefaultQueue().get_device()); - }); + // stable, so that devices the comparator sees as equal keep the order the platform reported + // them in and the device ids stay the same from run to run + std::stable_sort(this->availableDevices.begin(), + this->availableDevices.end(), + [](DeviceContext* c1, DeviceContext* c2) { + return compare(c1->queueBuffer.getDefaultQueue().get_device(), + c2->queueBuffer.getDefaultQueue().get_device()); + }); this->setDevice(0); this->deviceInitialized = true; @@ -73,8 +77,6 @@ void ConcreteAPI::finalize() { this->availableDevices.clear(); this->availableDevices.shrink_to_fit(); - this->graphs.clear(); - this->m_isFinalized = true; this->deviceInitialized = false; } @@ -89,8 +91,8 @@ int ConcreteAPI::getDeviceId() { } unsigned int ConcreteAPI::getGlobMemAlignment() { - auto device = this->currentDefaultQueue().get_device(); - return 128; // ToDo: find attribute; not: device.get_info(); + // ToDo: find attribute; not: device.get_info(); + return 128; } void ConcreteAPI::syncDevice() { this->currentQueueBuffer().syncAllQueuesWithHost(); } diff --git a/interfaces/sycl/DeviceCircularQueueBuffer.cpp b/interfaces/sycl/DeviceCircularQueueBuffer.cpp deleted file mode 100644 index c83eac2..0000000 --- a/interfaces/sycl/DeviceCircularQueueBuffer.cpp +++ /dev/null @@ -1,158 +0,0 @@ -// SPDX-FileCopyrightText: 2022 SeisSol Group -// -// SPDX-License-Identifier: BSD-3-Clause - -#include "DeviceCircularQueueBuffer.h" - -#include "Internals.h" -#include "SyclWrappedAPI.h" -#include "utils/logger.h" - -#include - -using namespace device::internals; - -namespace device { - -// very inconvenient, but AdaptiveCpp doesn't allow much freedom when constructing a property_list -#if defined(DEVICE_USE_GRAPH_CAPTURING) && defined(SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST) -#define BASE_QUEUE_PROPERTIES \ - sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::no_immediate_command_list {} -#else -#define BASE_QUEUE_PROPERTIES sycl::property::queue::in_order() -#endif - -QueueWrapper::QueueWrapper(const sycl::device& dev, - const std::function& handler) - : queue{dev, handler, sycl::property_list{BASE_QUEUE_PROPERTIES}} {} - -void QueueWrapper::synchronize() { waitCheck(queue); } -void QueueWrapper::dependency(QueueWrapper& other) { - // improvising... Adding an empty event here, mimicking a CUDA-like event dependency -#if defined(HIPSYCL_EXT_QUEUE_WAIT_LIST) || defined(ACPP_EXT_QUEUE_WAIT_LIST) || \ - defined(SYCL_EXT_ACPP_QUEUE_WAIT_LIST) - auto waitList1 = other.queue.get_wait_list(); - auto waitList2 = queue.get_wait_list(); - queue.submit([&](sycl::handler& h) { - h.depends_on(waitList1); - h.depends_on(waitList2); - DEVICE_SYCL_EMPTY_OPERATION(h); - }); -#else - auto queueEvent = other.queue.submit([&](sycl::handler& h) { DEVICE_SYCL_EMPTY_OPERATION(h); }); - queue.submit([&](sycl::handler& h) { DEVICE_SYCL_EMPTY_OPERATION_WITH_EVENT(h, queueEvent); }); -#endif -} - -DeviceCircularQueueBuffer::DeviceCircularQueueBuffer( - const sycl::device& dev, - const std::function& handler, - size_t capacity) - : queues{std::vector(capacity)}, deviceReference(dev), handlerReference(handler) { - if (capacity <= 0) - throw std::invalid_argument("Capacity must be at least 1!"); - - this->defaultQueue = QueueWrapper(dev, handler); - this->genericQueue = QueueWrapper(dev, handler); - for (size_t i = 0; i < capacity; i++) { - this->queues[i] = QueueWrapper(dev, handler); - } -} - -sycl::queue& DeviceCircularQueueBuffer::getDefaultQueue() { return defaultQueue.queue; } - -sycl::queue& DeviceCircularQueueBuffer::getGenericQueue() { return genericQueue.queue; } - -sycl::queue& DeviceCircularQueueBuffer::getNextQueue() { - (++this->counter) %= getCapacity(); - return (this->queues[this->counter].queue); -} - -std::vector DeviceCircularQueueBuffer::allQueues() { - std::vector queueCopy(queues.size() + 1); - queueCopy[0] = defaultQueue.queue; - for (size_t i = 0; i < queues.size(); ++i) { - queueCopy[i + 1] = queues[i].queue; - } - return queueCopy; -} - -sycl::queue* DeviceCircularQueueBuffer::newQueue(double priority) { - // missing for ACPP: how can we even find out the allowed priority range conveniently now? :/ - -#ifdef SYCL_EXT_ONEAPI_QUEUE_PRIORITY - const auto propertylist = [&]() -> sycl::property_list { - if (priority <= 0.33) { - return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_low()}; - } - if (priority >= 0.67) { - return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_high()}; - } - return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_normal()}; - }(); -#else - const auto propertylist{BASE_QUEUE_PROPERTIES}; -#endif - - auto* queue = new sycl::queue{deviceReference, handlerReference, propertylist}; - externalQueues.emplace_back(queue); - return queue; -} - -void DeviceCircularQueueBuffer::deleteQueue(void* queue) { - auto* queuePtr = static_cast(queue); - delete queuePtr; -} - -void DeviceCircularQueueBuffer::resetIndex() { this->counter = 0; } - -size_t DeviceCircularQueueBuffer::getCapacity() { return queues.size(); } - -void DeviceCircularQueueBuffer::forkQueueDepencency() { - for (auto& queue : this->queues) { - queue.dependency(defaultQueue); - } -} - -void DeviceCircularQueueBuffer::joinQueueDepencency() { - for (auto& queue : this->queues) { - defaultQueue.dependency(queue); - } -} - -void DeviceCircularQueueBuffer::syncQueueWithHost(sycl::queue* queuePtr) { waitCheck(*queuePtr); } - -void DeviceCircularQueueBuffer::syncAllQueuesWithHost() { - defaultQueue.synchronize(); - for (auto& q : this->queues) { - q.synchronize(); - } - for (auto* q : this->externalQueues) { - waitCheck(*q); - } -} - -bool DeviceCircularQueueBuffer::exists(sycl::queue* queuePtr) { - bool isDefaultQueue = queuePtr == (&defaultQueue.queue); - bool isGenericQueue = queuePtr == (&genericQueue.queue); - - bool isReservedQueue{true}; - for (auto& reservedQueue : queues) { - if (queuePtr != (&reservedQueue.queue)) { - isReservedQueue = false; - break; - } - } - - bool isExternalQueue = false; - for (auto& queue : externalQueues) { - if (queuePtr == queue) { - isExternalQueue = true; - break; - } - } - - return isDefaultQueue || isGenericQueue || isReservedQueue || isExternalQueue; -} - -} // namespace device diff --git a/interfaces/sycl/DeviceCircularQueueBuffer.h b/interfaces/sycl/DeviceCircularQueueBuffer.h deleted file mode 100644 index b24aace..0000000 --- a/interfaces/sycl/DeviceCircularQueueBuffer.h +++ /dev/null @@ -1,98 +0,0 @@ -// SPDX-FileCopyrightText: 2021 SeisSol Group -// -// SPDX-License-Identifier: BSD-3-Clause - -#ifndef SEISSOLDEVICE_INTERFACES_SYCL_DEVICECIRCULARQUEUEBUFFER_H_ -#define SEISSOLDEVICE_INTERFACES_SYCL_DEVICECIRCULARQUEUEBUFFER_H_ - -#include -#include -#include -#include -#include - -namespace device { -struct QueueWrapper { - sycl::queue queue; - QueueWrapper() = default; - QueueWrapper(const sycl::device& dev, const std::function& f); - - void synchronize(); - void dependency(QueueWrapper& other); -}; - -class DeviceCircularQueueBuffer { - public: - /* - * Creates a new circular buffer containing sycl queues. The buffer needs - * an async exception handler and a capacity that is currently per default 8. - */ - DeviceCircularQueueBuffer(const sycl::device& dev, - const std::function& f, - size_t capacity = 6); - - /* - * Returns the default queue created by a device. - */ - sycl::queue& getDefaultQueue(); - - /* - * Returns the generic queue created by a device. - * Note, this queue can be used for asynchronous - * memory copies - */ - sycl::queue& getGenericQueue(); - - /* - * Returns the next queue within the capacity of this buffer. - */ - sycl::queue& getNextQueue(); - - sycl::queue* newQueue(double priority); - - void deleteQueue(void* queue); - - std::vector allQueues(); - - /* - * Resets the index to the current element. - */ - void resetIndex(); - - /* - * Returns the default queue created by a device. - */ - size_t getCapacity(); - - /* - * Synchronizes a queue from the buffer with the host device. - */ - void syncQueueWithHost(sycl::queue* queuePtr); - - /* - * Synchronizes all queues from the buffer with the host device. - */ - void syncAllQueuesWithHost(); - - /* - *Returns true if the queue pointer is available on the buffer. - */ - bool exists(sycl::queue* queuePtr); - - void forkQueueDepencency(); - - void joinQueueDepencency(); - - private: - QueueWrapper defaultQueue; - QueueWrapper genericQueue; - std::vector queues; - std::vector externalQueues; - size_t counter; - sycl::device deviceReference; - std::function handlerReference; -}; - -} // namespace device - -#endif // SEISSOLDEVICE_INTERFACES_SYCL_DEVICECIRCULARQUEUEBUFFER_H_ diff --git a/interfaces/sycl/DeviceContext.cpp b/interfaces/sycl/DeviceContext.cpp index 1638390..f1568f9 100644 --- a/interfaces/sycl/DeviceContext.cpp +++ b/interfaces/sycl/DeviceContext.cpp @@ -8,9 +8,8 @@ #include "utils/logger.h" namespace device { -DeviceContext::DeviceContext(const sycl::device& targetDevice, size_t concurrencyLevel) - : queueBuffer{DeviceCircularQueueBuffer{ - targetDevice, [&](sycl::exception_list l) { onExceptionOccurred(l); }, concurrencyLevel}}, +DeviceContext::DeviceContext(const sycl::device& targetDevice) + : queueBuffer{targetDevice, [&](sycl::exception_list l) { onExceptionOccurred(l); }}, statistics{Statistics{}} {} void DeviceContext::onExceptionOccurred(sycl::exception_list& exceptions) { diff --git a/interfaces/sycl/DeviceContext.h b/interfaces/sycl/DeviceContext.h index 174cd44..72e602c 100644 --- a/interfaces/sycl/DeviceContext.h +++ b/interfaces/sycl/DeviceContext.h @@ -5,7 +5,7 @@ #ifndef SEISSOLDEVICE_INTERFACES_SYCL_DEVICECONTEXT_H_ #define SEISSOLDEVICE_INTERFACES_SYCL_DEVICECONTEXT_H_ -#include "DeviceCircularQueueBuffer.h" +#include "DeviceQueues.h" #include "Statistics.h" #include @@ -17,9 +17,9 @@ namespace device { */ class DeviceContext { public: - DeviceContext(const sycl::device& targetDevice, size_t concurrencyLevel); + explicit DeviceContext(const sycl::device& targetDevice); std::unordered_map memoryToSizeMap; - DeviceCircularQueueBuffer queueBuffer; + DeviceQueues queueBuffer; Statistics statistics; private: diff --git a/interfaces/sycl/DeviceQueues.cpp b/interfaces/sycl/DeviceQueues.cpp new file mode 100644 index 0000000..4046e0f --- /dev/null +++ b/interfaces/sycl/DeviceQueues.cpp @@ -0,0 +1,130 @@ +// SPDX-FileCopyrightText: 2022 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "DeviceQueues.h" + +#include "Internals.h" +#include "SyclWrappedAPI.h" +#include "utils/logger.h" + +#include +#include +#include +#include + +using namespace device::internals; + +namespace device { + +// very inconvenient, but AdaptiveCpp doesn't allow much freedom when constructing a property_list +#if defined(DEVICE_USE_GRAPH_CAPTURING) && defined(SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST) +#define IMMEDIATE_COMMAND_LIST_PROPERTY \ + , sycl::ext::intel::property::queue::no_immediate_command_list {} +#else +#define IMMEDIATE_COMMAND_LIST_PROPERTY +#endif + +// profiling is what makes the timings of AbstractAPI::timespanEvents available, and it costs on +// every submission, so it follows the build option rather than being on by default +#ifdef PROFILING_ENABLED +#define PROFILING_PROPERTY \ + , sycl::property::queue::enable_profiling {} +#else +#define PROFILING_PROPERTY +#endif + +#define BASE_QUEUE_PROPERTIES \ + sycl::property::queue::in_order {} \ + IMMEDIATE_COMMAND_LIST_PROPERTY PROFILING_PROPERTY + +DeviceQueues::DeviceQueues(const sycl::device& dev, + const std::function& handler) + : defaultQueue{dev, handler, sycl::property_list{BASE_QUEUE_PROPERTIES}}, deviceReference(dev), + handlerReference(handler) {} + +DeviceQueues::~DeviceQueues() { + if (!externalQueues.empty()) { + logInfo() << "DEVICE::WARNING:" << externalQueues.size() + << "device generic stream(s) were not deleted."; + } + for (auto* queue : externalQueues) { + delete queue; + } +} + +sycl::queue& DeviceQueues::getDefaultQueue() { return defaultQueue; } + +sycl::queue* DeviceQueues::newQueue(double priority) { + // missing for ACPP: how can we even find out the allowed priority range conveniently now? :/ + +#ifdef SYCL_EXT_ONEAPI_QUEUE_PRIORITY + const sycl::property_list propertylist = [&]() -> sycl::property_list { + if (std::isnan(priority)) { + return {BASE_QUEUE_PROPERTIES}; + } + if (priority <= 0.33) { + return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_low()}; + } + if (priority >= 0.67) { + return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_high()}; + } + return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_normal()}; + }(); +#else + const sycl::property_list propertylist{BASE_QUEUE_PROPERTIES}; +#endif + + auto* queue = new sycl::queue{deviceReference, handlerReference, propertylist}; + + const std::lock_guard lock(queueMutex); + externalQueues.emplace_back(queue); + return queue; +} + +void DeviceQueues::deleteQueue(void* queue) { + auto* queuePtr = static_cast(queue); + const std::lock_guard lock(queueMutex); + + // The queue has to leave the list before it is freed: syncAllQueuesWithHost walks that list, + // so a stale entry turns into a use-after-free at the next device-wide synchronization, far + // away from whoever destroyed the queue. + const auto entry = std::find(externalQueues.begin(), externalQueues.end(), queuePtr); + if (entry == externalQueues.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or belongs to a different device; not " + "freeing it again."; + return; + } + + externalQueues.erase(entry); + delete queuePtr; +} + +void DeviceQueues::syncQueueWithHost(sycl::queue* queuePtr) { waitCheck(*queuePtr); } + +void DeviceQueues::syncAllQueuesWithHost() { + waitCheck(defaultQueue); + + // copied under the lock: waiting on a queue takes as long as the work on it, and holding the + // lock for that would block every thread that wants to create or destroy one + const auto queues = [this]() { + const std::lock_guard lock(queueMutex); + return externalQueues; + }(); + + for (auto* queue : queues) { + waitCheck(*queue); + } +} + +bool DeviceQueues::exists(sycl::queue* queuePtr) { + if (queuePtr == &defaultQueue) { + return true; + } + + const std::lock_guard lock(queueMutex); + return std::find(externalQueues.begin(), externalQueues.end(), queuePtr) != externalQueues.end(); +} + +} // namespace device diff --git a/interfaces/sycl/DeviceQueues.h b/interfaces/sycl/DeviceQueues.h new file mode 100644 index 0000000..0566251 --- /dev/null +++ b/interfaces/sycl/DeviceQueues.h @@ -0,0 +1,69 @@ +// SPDX-FileCopyrightText: 2021 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#ifndef SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ +#define SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ + +#include +#include +#include +#include + +namespace device { + +/* + * Owns the queues of one device: the default queue every caller shares, and the queues handed out + * through newQueue. + */ +class DeviceQueues { + public: + DeviceQueues(const sycl::device& dev, const std::function& f); + ~DeviceQueues(); + + DeviceQueues(const DeviceQueues&) = delete; + DeviceQueues& operator=(const DeviceQueues&) = delete; + + /* + * Returns the default queue of this device. + */ + sycl::queue& getDefaultQueue(); + + /* + * Creates a queue owned by this device. `priority` follows the convention of + * AbstractAPI::createStream: 0 the lowest, 1 the highest, NaN the runtime default. + */ + sycl::queue* newQueue(double priority); + + /* + * Destroys a queue obtained from newQueue. + */ + void deleteQueue(void* queue); + + /* + * Synchronizes one queue with the host. + */ + void syncQueueWithHost(sycl::queue* queuePtr); + + /* + * Synchronizes every queue of this device with the host. + */ + void syncAllQueuesWithHost(); + + /* + * Returns true if the queue belongs to this device. + */ + bool exists(sycl::queue* queuePtr); + + private: + sycl::queue defaultQueue; + // guards externalQueues, which callers add to and remove from while other threads walk it + std::mutex queueMutex; + std::vector externalQueues; + sycl::device deviceReference; + std::function handlerReference; +}; + +} // namespace device + +#endif // SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ diff --git a/interfaces/sycl/DeviceType.cpp b/interfaces/sycl/DeviceType.cpp index 6a99a8a..d1737e9 100644 --- a/interfaces/sycl/DeviceType.cpp +++ b/interfaces/sycl/DeviceType.cpp @@ -7,6 +7,8 @@ #include "utils/env.h" #include +#include +#include namespace device { @@ -26,26 +28,21 @@ DeviceType fromSyclType(sycl::info::device_type type) { return DeviceType::OTHERS; } -bool compare(sycl::device devA, sycl::device devB) { - std::string env; - env += utils::Env("").get("PREFERRED_DEVICE_TYPE", ""); - - auto typeA = devA.get_info(); - auto typeB = devB.get_info(); - - if (convertToString(typeA).compare(env) == 0) { - return true; - } - if (convertToString(typeB).compare(env) == 0) { - return false; - } - - // devices of same type are sorted by their cl::deviceid - if (typeA == typeB) { - // return devA.get() < devB.get(); - } - - return fromSyclType(typeA) < fromSyclType(typeB); +bool compare(const sycl::device& devA, const sycl::device& devB) { + std::string preferred; + preferred += utils::Env("").get("PREFERRED_DEVICE_TYPE", ""); + + // A device of the preferred type sorts ahead of every other device, and the rest follow the + // order of the DeviceType enum. Deciding the two directions independently - as in "A wins if it + // matches, B wins if it matches" - makes both compare(a, b) and compare(b, a) true for two + // devices of the preferred type, and sorting on such a comparator is undefined. + const auto rank = [&preferred](const sycl::device& device) { + const auto type = device.get_info(); + const auto matchesPreference = convertToString(type) == preferred ? 0 : 1; + return std::make_pair(matchesPreference, static_cast(fromSyclType(type))); + }; + + return rank(devA) < rank(devB); } std::string convertToString(sycl::info::device_type type) { diff --git a/interfaces/sycl/DeviceType.h b/interfaces/sycl/DeviceType.h index 4a0e288..bf3243a 100644 --- a/interfaces/sycl/DeviceType.h +++ b/interfaces/sycl/DeviceType.h @@ -13,7 +13,7 @@ enum class DeviceType { GPU = 0, CPU = 1, FPGA = 2, HOST = 3, OTHERS = 4 }; std::string convertToString(sycl::info::device_type type); DeviceType fromSyclType(sycl::info::device_type type); -bool compare(sycl::device devA, sycl::device devB); +bool compare(const sycl::device& devA, const sycl::device& devB); } // namespace device diff --git a/interfaces/sycl/Events.cpp b/interfaces/sycl/Events.cpp index 27dfb0d..02691c1 100644 --- a/interfaces/sycl/Events.cpp +++ b/interfaces/sycl/Events.cpp @@ -19,15 +19,28 @@ using namespace device::internals; namespace { struct Event { std::optional syclEvent; + bool withTiming{false}; }; } // namespace -void* ConcreteAPI::createEvent(bool withTiming) { return new Event(); } +void* ConcreteAPI::createEvent(bool withTiming) { return new Event{std::nullopt, withTiming}; } double ConcreteAPI::timespanEvents(void* eventPtrStart, void* eventPtrEnd) { +#ifndef PROFILING_ENABLED + // Profiling information is only there when the queue carries the enable_profiling property, and + // asking an event without it throws from deep inside the runtime. + logError() << "The SYCL backend can only time events when it is built with " + "ENABLE_PROFILING_MARKERS=ON."; + return 0.0; +#else auto* start = static_cast(eventPtrStart); auto* end = static_cast(eventPtrEnd); + if (!(start->withTiming && end->withTiming)) { + logError() << "Timing was not requested for at least one of the events given for timing " + "calculation."; + } + if (!(start->syclEvent.has_value() && end->syclEvent.has_value())) { logError() << "Invalid events given for timing calculation."; } @@ -40,6 +53,7 @@ double ConcreteAPI::timespanEvents(void* eventPtrStart, void* eventPtrEnd) { // cf. https://oneapi-src.github.io/SYCLomatic/dev_guide/reference/diagnostic_ref/dpct1012.html return static_cast(endTime - startTime) / 1'000'000'000.0; +#endif } void ConcreteAPI::destroyEvent(void* eventPtr) { diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index 1905a6a..8122501 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -7,25 +7,53 @@ #include "SyclWrappedAPI.h" #include "utils/logger.h" +#include #include +#include +#include #include using namespace device; -/* This is a wrapped graph capturing CUDA mechanism. - * Call the following in order to capture a computational graph - * streamBeginCapture(); // 1 +/* Two ways of building a compute graph are offered; see the CUDA backend for the shapes. * - * // your GPU code here // 2 + * Both rest on the same mechanism here: the oneAPI graph extension records queues. Whole-queue + * recording captures everything submitted between begin and end. Node construction records one + * segment per node and expresses the edges through barriers on the recorded events, which is + * what the extension offers in place of the node handles the CUDA and HIP backends hand out. * - * streamEndCapture(); // 3 - * auto graph = getGraphInstance(); // 4 - * - * Once you have a compute-graph recorded you can invoke it as follows: - * launchGraph(graph) // 1 - * syncGraph(graph) // 2 + * The queues record in order, so submissions within one node are chained automatically, and two + * nodes recorded onto the same queue end up ordered even without an edge between them. Siblings + * that are meant to run concurrently therefore have to be recorded onto different queues. * */ +namespace device { +struct DeviceGraph { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + std::optional> + instance; + sycl::ext::oneapi::experimental::command_graph< + sycl::ext::oneapi::experimental::graph_state::modifiable> + graph; + + // one entry per graphEndNode call; SYCL expresses graph edges through the events of recorded + // submissions rather than through node objects + std::vector> nodes; + + // queues that recording has been started on, so that it is only started once per queue + std::vector recordedQueues; + + DeviceGraph(const sycl::context& context, const sycl::device& device) : graph(context, device) {} +#endif + + bool ready{false}; + + DeviceGraph(const DeviceGraph&) = delete; + DeviceGraph& operator=(const DeviceGraph&) = delete; +}; +} // namespace device + bool ConcreteAPI::isCapableOfGraphCapturing() { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT return true; @@ -34,52 +62,134 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { #endif } -DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { - auto handle = DeviceGraphHandle(); +bool ConcreteAPI::isCapableOfGraphNodes() { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - std::vector queues; + return true; +#else + return false; +#endif +} + +DeviceGraphHandle ConcreteAPI::streamBeginCapture(const std::vector& streamPtrs) { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + if (streamPtrs.empty()) { + logError() << "Graph capturing records queues, so it needs at least one."; + return DeviceGraphHandle(); + } + std::vector queues; + queues.reserve(streamPtrs.size()); for (auto* streamPtr : streamPtrs) { queues.emplace_back(*static_cast(streamPtr)); } - auto recordingGraph = sycl::ext::oneapi::experimental::command_graph< - sycl::ext::oneapi::experimental::graph_state::modifiable>(queues.at(0).get_context(), - queues.at(0).get_device()); + auto graphInstance = + std::make_shared(queues.at(0).get_context(), queues.at(0).get_device()); + graphInstance->graph.begin_recording(queues); + + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); +#endif +} + +void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto* graphInstance = handle.get(); + assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + assert(!graphInstance->instance.has_value() && "a graph is instantiated once"); + + graphInstance->graph.end_recording(); + graphInstance->instance = std::optional>(graphInstance->graph.finalize()); + + graphInstance->ready = true; +#endif +} + +DeviceGraphHandle ConcreteAPI::graphCreate() { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto& queue = this->currentDefaultQueue(); + return DeviceGraphHandle(std::make_shared(queue.get_context(), queue.get_device())); +#else + return DeviceGraphHandle(); +#endif +} + +void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); + assert(!graphInstance->ready && "no nodes can be added to an instantiated graph"); - { - std::lock_guard guard(apiMutex); - graphs.push_back(GraphDetails{std::nullopt, std::move(recordingGraph), false}); - handle = DeviceGraphHandle(graphs.size() - 1); + auto* queue = static_cast(streamPtr); + auto& recorded = graphInstance->recordedQueues; + if (std::find(recorded.begin(), recorded.end(), queue) == recorded.end()) { + graphInstance->graph.begin_recording(*queue); + recorded.push_back(queue); + } - GraphDetails& graphInstance = graphs[handle.getGraphId()]; + std::vector nativeDependencies; + for (const auto& dependency : dependencies) { + assert(dependency.isInitialized() && "an uninitialized node cannot be depended upon"); + const auto& events = graphInstance->nodes.at(dependency.getNodeId()); + nativeDependencies.insert(nativeDependencies.end(), events.begin(), events.end()); + } - graphInstance.graph.begin_recording(queues); + if (!nativeDependencies.empty()) { + // an empty node that pulls the dependencies onto this queue. Everything the caller records + // next follows it, because the queues are in order. + queue->submit([&nativeDependencies](sycl::handler& handler) { + handler.ext_oneapi_barrier(nativeDependencies); + }); } #endif - return handle; } -void ConcreteAPI::streamEndCapture(DeviceGraphHandle handle) { +DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a node must be opened before it can be closed"); + + auto* queue = static_cast(streamPtr); + + // Closing through a command group rather than through queue::ext_oneapi_submit_barrier: on an + // in-order queue the shortcut hands back the last recorded event, which is not the barrier's + // own event if the preceding submission discarded its event. + std::vector produced{ + queue->submit([](sycl::handler& handler) { handler.ext_oneapi_barrier(); })}; + + graphInstance->nodes.emplace_back(std::move(produced)); + return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); +#else + return DeviceGraphNodeHandle(); +#endif +} + +void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - std::lock_guard guard(apiMutex); - auto& graphInstance = graphs[handle.getGraphId()]; - graphInstance.graph.end_recording(); - graphInstance.instance = std::optional>(graphInstance.graph.finalize()); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + assert(!graphInstance->instance.has_value() && "a graph is instantiated once"); + + graphInstance->graph.end_recording(); + graphInstance->instance = std::optional>(graphInstance->graph.finalize()); - graphInstance.ready = true; + graphInstance->ready = true; #endif } -void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) { +void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - assert(graphHandle.isInitialized() && "a graph must be captured before launching"); - GraphDetails graphInstance = [&]() { - std::lock_guard guard(apiMutex); - return graphs[graphHandle.getGraphId()]; - }(); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && graphInstance->ready && + "a graph must be captured before launching"); + static_cast(streamPtr)->submit( - [&](sycl::handler& handler) { handler.ext_oneapi_graph(graphInstance.instance.value()); }); + [&](sycl::handler& handler) { handler.ext_oneapi_graph(graphInstance->instance.value()); }); #endif } diff --git a/interfaces/sycl/Memory.cpp b/interfaces/sycl/Memory.cpp index 860365c..cf6e0f5 100644 --- a/interfaces/sycl/Memory.cpp +++ b/interfaces/sycl/Memory.cpp @@ -6,36 +6,40 @@ #include "SyclWrappedAPI.h" #include +#include using namespace device; using namespace device::internals; void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { + const std::lock_guard lock(apiMutex); + auto* ptr = malloc_device(size, this->currentDefaultQueue()); this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); - waitCheck(this->currentDefaultQueue()); return ptr; } void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) { + const std::lock_guard lock(apiMutex); + auto* ptr = malloc_shared(size, this->currentDefaultQueue()); this->currentStatistics().allocatedUnifiedMemBytes += size; this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); - waitCheck(this->currentDefaultQueue()); return ptr; } void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) { + const std::lock_guard lock(apiMutex); + auto* ptr = malloc_host(size, this->currentDefaultQueue()); this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); - waitCheck(this->currentDefaultQueue()); return ptr; } -void ConcreteAPI::freeMem(void* devPtr) { +void ConcreteAPI::freeMem(void* devPtr, bool unified) { // NOTE: Freeing nullptr results in segfault in oneAPI. It is an opposite behaviour // contrast to C++/CUDA/HIP if (devPtr == nullptr) { @@ -50,6 +54,8 @@ void ConcreteAPI::freeMem(void* devPtr) { return; } + const std::lock_guard lock(apiMutex); + // Use the first device context to free memory DeviceContext* context = this->availableDevices[getDeviceId()]; if (!context) { @@ -61,11 +67,17 @@ void ConcreteAPI::freeMem(void* devPtr) { return; // the std::throw is throwing some errors during the program finalization } - context->statistics.deallocatedMemBytes += map.at(devPtr); + const auto size = map.at(devPtr); + context->statistics.deallocatedMemBytes += size; + if (unified) { + context->statistics.allocatedUnifiedMemBytes -= size; + } map.erase(devPtr); + // freeing memory that a queue may still be reading from is undefined, and the caller has no + // way to state that it is done, so the wait stays auto& queue = context->queueBuffer.getDefaultQueue(); - sycl::free(devPtr, queue.get_context()); queue.wait(); + sycl::free(devPtr, queue.get_context()); } void ConcreteAPI::freeGlobMem(void* devPtr) { @@ -80,7 +92,7 @@ void ConcreteAPI::freeUnifiedMem(void* devPtr) { // NOTE: Freeing nullptr results in segfault in oneAPI. It is an opposite behavior // contrast to C++/CUDA/HIP if (devPtr != nullptr) { - this->freeMem(devPtr); + this->freeMem(devPtr, true); } } @@ -106,6 +118,8 @@ void ConcreteAPI::freeMemAsync(void* devPtr, void* streamPtr) { } std::string ConcreteAPI::getMemLeaksReport() { + const std::lock_guard lock(apiMutex); + std::ostringstream report{}; report << "----MEMORY REPORT----\n"; @@ -124,10 +138,14 @@ size_t ConcreteAPI::getMaxAvailableMem() { } size_t ConcreteAPI::getCurrentlyOccupiedMem() { + const std::lock_guard lock(apiMutex); + return this->currentStatistics().allocatedMemBytes; } size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { + const std::lock_guard lock(apiMutex); + return this->currentStatistics().allocatedUnifiedMemBytes; } diff --git a/interfaces/sycl/Streams.cpp b/interfaces/sycl/Streams.cpp index 3e96c24..aaa6a32 100644 --- a/interfaces/sycl/Streams.cpp +++ b/interfaces/sycl/Streams.cpp @@ -43,8 +43,16 @@ bool ConcreteAPI::isStreamWorkDone(void* streamPtr) { // otherwise, synchronize #ifdef SYCL_EXT_ONEAPI_QUEUE_EMPTY return queuePtr->ext_oneapi_empty(); -#elif defined(HIPSYCL_EXT_QUEUE_WAIT_LIST) || defined(ACPP_EXT_QUEUE_WAIT_LIST) - return queuePtr->get_wait_list().empty(); +#elif defined(HIPSYCL_EXT_QUEUE_WAIT_LIST) || defined(ACPP_EXT_QUEUE_WAIT_LIST) || \ + defined(SYCL_EXT_ACPP_QUEUE_WAIT_LIST) + // The wait list holds the events a newly submitted operation would have to depend on. Those + // entries are not dropped once they have been reached, so an empty list means "nothing was + // ever submitted", not "nothing is outstanding". Ask the events themselves instead. + const auto waitList = queuePtr->get_wait_list(); + return std::all_of(waitList.begin(), waitList.end(), [](const sycl::event& event) { + return event.get_info() == + sycl::info::event_command_status::complete; + }); #else this->currentQueueBuffer().syncQueueWithHost(queuePtr); return true; @@ -67,7 +75,9 @@ void ConcreteAPI::streamWaitMemory(void* streamPtr, uint32_t* location, uint32_t volatile uint32_t* spinLocation = location; queuePtr->single_task([=]() { while (true) { - if (*spinLocation == value) { + // ">=", like the wait-value operations of the other backends: a counter that is written + // once per step is past the awaited value by the time this runs often enough + if (*spinLocation >= value) { return; } diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index 3f854b4..2d9ed02 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -42,7 +42,7 @@ #define DEVICE_SYCL_EMPTY_OPERATION_WITH_EVENT(handle, event) \ handle.depends_on(event); \ handle.DEVICE_SYCL_DIRECT_OPERATION_NAME([=](...) {}); -#elif defined(SYCL_EXT_ONEAPI_ENQUEUE_BARRIER) && !defined(DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT) +#elif defined(SYCL_EXT_ONEAPI_ENQUEUE_BARRIER) #define DEVICE_SYCL_EMPTY_OPERATION(handle) handle.ext_oneapi_barrier(); #define DEVICE_SYCL_EMPTY_OPERATION_WITH_EVENT(handle, event) handle.ext_oneapi_barrier({event}); #else @@ -117,9 +117,18 @@ class ConcreteAPI : public AbstractAPI { void syncDefaultStreamWithHost() override; bool isCapableOfGraphCapturing() override; - DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; - void streamEndCapture(DeviceGraphHandle handle) override; - void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override; + DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) override; + void streamEndCapture(const DeviceGraphHandle& handle) override; + void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; + + bool isCapableOfGraphNodes() override; + DeviceGraphHandle graphCreate() override; + void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) override; + DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) override; + void graphInstantiate(const DeviceGraphHandle& graphHandle) override; void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; @@ -153,31 +162,13 @@ class ConcreteAPI : public AbstractAPI { DeviceContext* currentContext() { return this->availableDevices[getDeviceId()]; } sycl::queue& currentDefaultQueue() { return this->currentQueueBuffer().getDefaultQueue(); } - DeviceCircularQueueBuffer& currentQueueBuffer() { return this->currentContext()->queueBuffer; } + DeviceQueues& currentQueueBuffer() { return this->currentContext()->queueBuffer; } Statistics& currentStatistics() { return this->currentContext()->statistics; } std::unordered_map& currentMemoryToSizeMap() { return this->currentContext()->memoryToSizeMap; } -#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - struct GraphDetails { - std::optional> - instance; - sycl::ext::oneapi::experimental::command_graph< - sycl::ext::oneapi::experimental::graph_state::modifiable> - graph; - bool ready{false}; - }; -#else - struct GraphDetails { - bool ready{false}; - }; -#endif - - std::vector graphs; - - void freeMem(void* devPtr); + void freeMem(void* devPtr, bool unified = false); void initDevices(); diff --git a/sycl.cmake b/sycl.cmake index 666e1e7..8b2ff28 100644 --- a/sycl.cmake +++ b/sycl.cmake @@ -10,7 +10,7 @@ set(DEVICE_SOURCE_FILES device.cpp interfaces/sycl/Memory.cpp interfaces/sycl/Streams.cpp interfaces/sycl/DeviceContext.cpp - interfaces/sycl/DeviceCircularQueueBuffer.cpp + interfaces/sycl/DeviceQueues.cpp interfaces/sycl/DeviceType.cpp algorithms/sycl/ArrayManip.cpp algorithms/sycl/BatchManip.cpp diff --git a/tests/BaseTestSuite.h b/tests/BaseTestSuite.h index 4d2c2ae..7583a94 100644 --- a/tests/BaseTestSuite.h +++ b/tests/BaseTestSuite.h @@ -14,18 +14,13 @@ using namespace device; using namespace ::testing; -static bool setUp = false; - class BaseTestSuite : public ::testing::Test { public: - DeviceInstance* device; + DeviceInstance* device{nullptr}; BaseTestSuite() { randomEngine.seed(randomDevice()); } - void SetUp() { - device = &DeviceInstance::getInstance(); - setUp = true; - } + void SetUp() override { device = &DeviceInstance::getInstance(); } protected: std::random_device randomDevice; diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 7369f47..d4607d3 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -8,8 +8,8 @@ project(examples) #set(CMAKE_CXX_CLANG_TIDY clang-tidy) set(SM "sm_60" CACHE STRING "model of Nvidia Streaming Multiprocessor") -set(SM_OPTIONS "sm_50" "sm_60" "sm_70" "sm_71" "sm_75" "sm_80" "sm_86" - "gfx906" "gfx908 " +set(SM_OPTIONS "sm_50" "sm_60" "sm_70" "sm_71" "sm_75" "sm_80" "sm_86" "sm_89" "sm_90" + "gfx906" "gfx908" "gfx90a" "gfx942" "dg1" "bdw" "skl" "Gen8" "Gen9" "Gen11" "Gen12LP") set_property(CACHE SM PROPERTY STRINGS ${SM_OPTIONS}) @@ -22,7 +22,15 @@ add_subdirectory(.. root) find_package(GTest REQUIRED) -add_executable(tests main.cpp reductions.cpp memory.cpp array_manip.cpp batch_manip.cpp) +add_executable(tests + main.cpp + reductions.cpp + memory.cpp + array_manip.cpp + batch_manip.cpp + batch_transfer.cpp + streams.cpp + graphs.cpp) target_link_libraries(tests PRIVATE device ${GTEST_BOTH_LIBRARIES}) target_include_directories(tests PRIVATE ${GTEST_INCLUDE_DIR}) diff --git a/tests/array_manip.cpp b/tests/array_manip.cpp index 71df856..18e30fd 100644 --- a/tests/array_manip.cpp +++ b/tests/array_manip.cpp @@ -131,3 +131,16 @@ TEST_F(ArrayManip, scale) { device->api->freeGlobMem(arr); } + +TEST_F(ArrayManip, anEmptyArrayIsNoWork) { + auto* devArray = static_cast(device->api->allocGlobMem(sizeof(float))); + + device->algorithms.fillArray(devArray, 1.0F, 0, device->api->getDefaultStream()); + device->algorithms.scaleArray(devArray, 2.0F, 0, device->api->getDefaultStream()); + device->algorithms.touchMemory(devArray, 0, true, device->api->getDefaultStream()); + device->api->syncDefaultStreamWithHost(); + + SUCCEED(); + + device->api->freeGlobMem(devArray); +} diff --git a/tests/batch_transfer.cpp b/tests/batch_transfer.cpp new file mode 100644 index 0000000..b632f30 --- /dev/null +++ b/tests/batch_transfer.cpp @@ -0,0 +1,228 @@ +// SPDX-FileCopyrightText: 2026 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "BaseTestSuite.h" +#include "device.h" + +#include "gtest/gtest.h" +#include +#include + +using namespace device; +using namespace ::testing; + +namespace { +constexpr std::size_t BatchSize = 64; +constexpr std::size_t ElementSize = 48; +} // namespace + +/** + * Covers the batched transfers and the pointer arithmetic helper, which the other suites do not + * touch. These are the operations the solver uses to move data in and out of batched buffers, so + * a wrong stride or a skipped entry shows up as a wrong answer far away from here. + */ +class BatchTransfer : public BaseTestSuite { + public: + void SetUp() override { + BaseTestSuite::SetUp(); + stream = device->api->createStream(); + src = static_cast(device->api->allocGlobMem(BatchSize * ElementSize * sizeof(float))); + dst = static_cast(device->api->allocGlobMem(BatchSize * ElementSize * sizeof(float))); + srcBatch = static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + dstBatch = static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + + for (std::size_t i = 0; i < BatchSize; ++i) { + srcBatch[i] = src + i * ElementSize; + dstBatch[i] = dst + i * ElementSize; + } + } + + void TearDown() override { + device->api->freeUnifiedMem(dstBatch); + device->api->freeUnifiedMem(srcBatch); + device->api->freeGlobMem(dst); + device->api->freeGlobMem(src); + device->api->destroyGenericStream(stream); + } + + protected: + void upload(float* target, const std::vector& host) { + device->api->copyToAsync(target, host.data(), host.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + } + + std::vector download(const float* source) { + std::vector host(BatchSize * ElementSize, -1); + device->api->copyFromAsync(host.data(), source, host.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + return host; + } + + // a value that differs per batch entry and per element, so a mixed-up stride cannot pass + static float pattern(std::size_t entry, std::size_t element) { + return static_cast(entry * ElementSize + element); + } + + void* stream{nullptr}; + float* src{nullptr}; + float* dst{nullptr}; + float** srcBatch{nullptr}; + float** dstBatch{nullptr}; +}; + +TEST_F(BatchTransfer, streamBatchedDataCopiesEveryEntry) { + std::vector hostSrc(BatchSize * ElementSize); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + hostSrc[i * ElementSize + j] = pattern(i, j); + } + } + upload(src, hostSrc); + upload(dst, std::vector(BatchSize * ElementSize, 0.0F)); + + device->algorithms.streamBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + const auto hostDst = download(dst); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + ASSERT_EQ(pattern(i, j), hostDst[i * ElementSize + j]) + << "at entry " << i << ", element " << j; + } + } +} + +TEST_F(BatchTransfer, streamBatchedDataSkipsNullEntries) { + upload(src, std::vector(BatchSize * ElementSize, 1.0F)); + upload(dst, std::vector(BatchSize * ElementSize, -3.0F)); + + for (std::size_t i = 0; i < BatchSize; i += 2) { + srcBatch[i] = nullptr; + } + + device->algorithms.streamBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + const auto hostDst = download(dst); + for (std::size_t i = 0; i < BatchSize; ++i) { + const float expected = (i % 2 == 0) ? -3.0F : 1.0F; + for (std::size_t j = 0; j < ElementSize; ++j) { + ASSERT_EQ(expected, hostDst[i * ElementSize + j]) << "at entry " << i << ", element " << j; + } + } +} + +TEST_F(BatchTransfer, accumulateBatchedDataAdds) { + std::vector hostSrc(BatchSize * ElementSize); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + hostSrc[i * ElementSize + j] = pattern(i, j); + } + } + upload(src, hostSrc); + upload(dst, std::vector(BatchSize * ElementSize, 5.0F)); + + device->algorithms.accumulateBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + const auto hostDst = download(dst); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + ASSERT_EQ(5.0F + pattern(i, j), hostDst[i * ElementSize + j]) + << "at entry " << i << ", element " << j; + } + } +} + +TEST_F(BatchTransfer, accumulateBatchedDataIsRepeatable) { + upload(src, std::vector(BatchSize * ElementSize, 2.0F)); + upload(dst, std::vector(BatchSize * ElementSize, 0.0F)); + + for (int round = 0; round < 3; ++round) { + device->algorithms.accumulateBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + } + device->api->syncStreamWithHost(stream); + + for (const auto value : download(dst)) { + ASSERT_EQ(6.0F, value); + } +} + +TEST_F(BatchTransfer, incrementalAddBuildsAStridedPointerTable) { + auto** table = static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + + device->algorithms.incrementalAdd(table, src, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + // the stride is given in elements, not bytes + for (std::size_t i = 0; i < BatchSize; ++i) { + ASSERT_EQ(src + i * ElementSize, table[i]) << "at entry " << i; + } + + device->api->freeUnifiedMem(table); +} + +/** + * The copy routines step down from 16-byte accesses, and an element stride that is not a multiple + * of 16 bytes puts every second element off that boundary. 47 floats are 188 bytes, so entry 1 + * starts 4-byte aligned and a 16-byte access to it faults. + */ +TEST_F(BatchTransfer, unalignedElementsAreCopied) { + constexpr std::size_t OddElementSize = 47; + + auto* oddSrc = + static_cast(device->api->allocGlobMem(BatchSize * OddElementSize * sizeof(float))); + auto* oddDst = + static_cast(device->api->allocGlobMem(BatchSize * OddElementSize * sizeof(float))); + auto** oddSrcBatch = + static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + auto** oddDstBatch = + static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + + std::vector hostSrc(BatchSize * OddElementSize); + for (std::size_t i = 0; i < BatchSize; ++i) { + oddSrcBatch[i] = oddSrc + i * OddElementSize; + oddDstBatch[i] = oddDst + i * OddElementSize; + for (std::size_t j = 0; j < OddElementSize; ++j) { + hostSrc[i * OddElementSize + j] = static_cast(i * OddElementSize + j); + } + } + + const std::vector zeroes(hostSrc.size(), 0.0F); + device->api->copyToAsync(oddSrc, hostSrc.data(), hostSrc.size() * sizeof(float), stream); + device->api->copyToAsync(oddDst, zeroes.data(), zeroes.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + + device->algorithms.streamBatchedData( + const_cast(oddSrcBatch), oddDstBatch, OddElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + std::vector hostDst(hostSrc.size(), -1); + device->api->copyFromAsync(hostDst.data(), oddDst, hostDst.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + + for (std::size_t i = 0; i < hostSrc.size(); ++i) { + ASSERT_EQ(hostSrc[i], hostDst[i]) << "at " << i; + } + + device->api->freeUnifiedMem(oddDstBatch); + device->api->freeUnifiedMem(oddSrcBatch); + device->api->freeGlobMem(oddDst); + device->api->freeGlobMem(oddSrc); +} + +TEST_F(BatchTransfer, anEmptyBatchIsNoWork) { + device->algorithms.streamBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, 0, stream); + device->algorithms.accumulateBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, 0, stream); + device->algorithms.touchBatchedMemory(dstBatch, ElementSize, 0, true, stream); + device->api->syncStreamWithHost(stream); + + SUCCEED(); +} diff --git a/tests/graphs.cpp b/tests/graphs.cpp new file mode 100644 index 0000000..cb161f5 --- /dev/null +++ b/tests/graphs.cpp @@ -0,0 +1,369 @@ +// SPDX-FileCopyrightText: 2026 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "BaseTestSuite.h" +#include "device.h" + +#include "gtest/gtest.h" +#include +#include +#include + +using namespace device; +using namespace ::testing; + +namespace { +constexpr std::size_t ArraySize = 1 << 14; +constexpr std::size_t BranchCount = 4; +constexpr std::size_t ChunkSize = ArraySize / BranchCount; +} // namespace + +/** + * The tests below check what a graph guarantees, not how it is built: every one of them states a + * dependency structure and then asserts a result that only comes out right if that structure was + * honoured. Operations are picked so that swapping two of them changes the answer - a fill after a + * scale does not give the same array as a scale after a fill - because operations that commute + * would pass no matter how the edges came out. + */ +class Graphs : public BaseTestSuite { + public: + void SetUp() override { + BaseTestSuite::SetUp(); + devArray = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); + mainStream = device->api->createStream(); + for (auto& stream : branchStreams) { + stream = device->api->createStream(); + } + } + + void TearDown() override { + for (auto* stream : branchStreams) { + device->api->destroyGenericStream(stream); + } + device->api->destroyGenericStream(mainStream); + device->api->freeGlobMem(devArray); + } + + protected: + std::vector download() { + std::vector host(ArraySize, -1); + device->api->copyFromAsync(host.data(), devArray, ArraySize * sizeof(float), mainStream); + device->api->syncStreamWithHost(mainStream); + return host; + } + + void fill(float value) { + device->algorithms.fillArray(devArray, value, ArraySize, mainStream); + device->api->syncStreamWithHost(mainStream); + } + + static void expectChunk(const std::vector& host, std::size_t chunk, float value) { + for (std::size_t i = chunk * ChunkSize; i < (chunk + 1) * ChunkSize; ++i) { + ASSERT_EQ(value, host[i]) << "at index " << i << " of chunk " << chunk; + } + } + + bool graphCapturingUnavailable() { return !device->api->isCapableOfGraphCapturing(); } + + bool graphNodesUnavailable() { return !device->api->isCapableOfGraphNodes(); } + + float* devArray{nullptr}; + void* mainStream{nullptr}; + std::array branchStreams{}; +}; + +TEST_F(Graphs, captureReplaysTheRecordedSequence) { + if (graphCapturingUnavailable()) { + GTEST_SKIP() << "the backend does not support graph capturing"; + } + + fill(0); + + std::vector streams{mainStream}; + auto graph = device->api->streamBeginCapture(streams); + device->algorithms.fillArray(devArray, 1.0F, ArraySize, mainStream); + device->algorithms.scaleArray(devArray, 2.0F, ArraySize, mainStream); + device->api->streamEndCapture(graph); + + ASSERT_TRUE(graph.isInitialized()); + + // nothing has run yet: capturing records, it does not execute + for (const auto value : download()) { + ASSERT_EQ(0.0F, value); + } + + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(2.0F, value); + } + + // the fill is part of the graph, so a second replay lands on the same value rather than doubling + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(2.0F, value); + } +} + +TEST_F(Graphs, captureRecordsCrossStreamEvents) { + if (graphCapturingUnavailable()) { + GTEST_SKIP() << "the backend does not support graph capturing"; + } + + fill(0); + + // the fork/join shape that the stream path uses: an event hands work from the recorded stream + // to a side stream and back. Inside a capture these become graph edges rather than real waits. + auto* forkEvent = device->api->createEvent(); + auto* joinEvent = device->api->createEvent(); + + std::vector streams{mainStream, branchStreams[0]}; + auto graph = device->api->streamBeginCapture(streams); + + device->algorithms.fillArray(devArray, 3.0F, ArraySize, mainStream); + device->api->recordEventOnStream(forkEvent, mainStream); + device->api->syncStreamWithEvent(branchStreams[0], forkEvent); + device->algorithms.scaleArray(devArray, 4.0F, ArraySize, branchStreams[0]); + device->api->recordEventOnStream(joinEvent, branchStreams[0]); + device->api->syncStreamWithEvent(mainStream, joinEvent); + + device->api->streamEndCapture(graph); + ASSERT_TRUE(graph.isInitialized()); + + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(12.0F, value); + } + + device->api->destroyEvent(joinEvent); + device->api->destroyEvent(forkEvent); +} + +TEST_F(Graphs, nodesRunInDependencyOrder) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(0); + + auto graph = device->api->graphCreate(); + ASSERT_TRUE(graph.isInitialized()); + + const auto first = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 1.0F, ArraySize, stream); + }); + const auto second = device->api->graphAddNode(graph, {first}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 3.0F, ArraySize, stream); + }); + const auto third = device->api->graphAddNode(graph, {second}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 5.0F, ArraySize, stream); + }); + ASSERT_TRUE(third.isInitialized()); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + // the fill has to come first: if it ran last, every entry would be 1 instead + for (const auto value : download()) { + ASSERT_EQ(15.0F, value); + } +} + +TEST_F(Graphs, nodesForkAndJoin) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(-1); + + auto graph = device->api->graphCreate(); + + const auto root = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 0.0F, ArraySize, stream); + }); + + // each branch owns a disjoint chunk and runs on its own stream, so they may overlap + std::vector branches; + for (std::size_t i = 0; i < BranchCount; ++i) { + branches.push_back( + device->api->graphAddNode(graph, {root}, branchStreams[i], [&, i](void* stream) { + device->algorithms.fillArray( + devArray + i * ChunkSize, static_cast(i + 1), ChunkSize, stream); + })); + } + + const auto join = device->api->graphAddNode(graph, branches, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 10.0F, ArraySize, stream); + }); + ASSERT_TRUE(join.isInitialized()); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + // a chunk holding i+1 means the join overtook its branch; a chunk holding 0 means the root + // overtook it + const auto host = download(); + for (std::size_t i = 0; i < BranchCount; ++i) { + expectChunk(host, i, 10.0F * static_cast(i + 1)); + } +} + +TEST_F(Graphs, anEmptyNodeJoinsItsDependencies) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(-1); + + auto graph = device->api->graphCreate(); + + const auto root = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 0.0F, ArraySize, stream); + }); + + std::vector branches; + for (std::size_t i = 0; i < BranchCount; ++i) { + branches.push_back( + device->api->graphAddNode(graph, {root}, branchStreams[i], [&, i](void* stream) { + device->algorithms.fillArray( + devArray + i * ChunkSize, static_cast(i + 1), ChunkSize, stream); + })); + } + + // a node that records nothing stands for its own dependencies, which is what makes it usable + // as a join without costing a command + const auto join = device->api->graphAddNode(graph, branches, mainStream, [](void*) {}); + + const auto last = device->api->graphAddNode(graph, {join}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 100.0F, ArraySize, stream); + }); + ASSERT_TRUE(last.isInitialized()); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + const auto host = download(); + for (std::size_t i = 0; i < BranchCount; ++i) { + expectChunk(host, i, 100.0F * static_cast(i + 1)); + } +} + +TEST_F(Graphs, aNodeGraphCanBeLaunchedRepeatedly) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(1); + + // no fill inside the graph, so repeated launches accumulate and a graph that silently ran only + // once would be caught + auto graph = device->api->graphCreate(); + device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 2.0F, ArraySize, stream); + }); + device->api->graphInstantiate(graph); + + for (int i = 0; i < 3; ++i) { + device->api->launchGraph(graph, mainStream); + } + device->api->syncStreamWithHost(mainStream); + + for (const auto value : download()) { + ASSERT_EQ(8.0F, value); + } +} + +TEST_F(Graphs, handlesOwnTheirGraph) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + DeviceGraphHandle empty; + EXPECT_FALSE(empty.isInitialized()); + EXPECT_TRUE(!empty); + + auto graph = device->api->graphCreate(); + device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 5.0F, ArraySize, stream); + }); + device->api->graphInstantiate(graph); + + auto copy = graph; + EXPECT_TRUE(copy.isInitialized()); + graph.reset(); + EXPECT_FALSE(graph.isInitialized()); + + // the graph is still alive through the second handle + device->api->launchGraph(copy, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(5.0F, value); + } +} + +TEST_F(Graphs, droppedGraphsReleaseTheirResources) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + // A workload that keys its graphs on something that varies - a time step width, say - builds + // and drops them all the time, and that has to stay flat rather than accumulate device-side + // resources. + for (int i = 0; i < 256; ++i) { + auto graph = device->api->graphCreate(); + device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, static_cast(i), ArraySize, stream); + }); + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + } + device->api->syncStreamWithHost(mainStream); + + for (const auto value : download()) { + ASSERT_EQ(255.0F, value); + } +} + +TEST_F(Graphs, siblingNodesMayShareAStream) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(-1); + + // Two nodes with the same dependency and no edge between them, recorded onto one stream. They + // may end up ordered - a backend that expresses edges through the recorded stream orders them - + // but they write disjoint chunks, so the result is the same either way and neither may be lost. + auto graph = device->api->graphCreate(); + + const auto root = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 0.0F, ArraySize, stream); + }); + + std::vector siblings; + for (std::size_t i = 0; i < BranchCount; ++i) { + siblings.push_back(device->api->graphAddNode(graph, {root}, mainStream, [&, i](void* stream) { + device->algorithms.fillArray( + devArray + i * ChunkSize, static_cast(i + 1), ChunkSize, stream); + })); + } + + device->api->graphAddNode(graph, siblings, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 10.0F, ArraySize, stream); + }); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + const auto host = download(); + for (std::size_t i = 0; i < BranchCount; ++i) { + expectChunk(host, i, 10.0F * static_cast(i + 1)); + } +} diff --git a/tests/reductions.cpp b/tests/reductions.cpp index 0aaac48..00b9243 100644 --- a/tests/reductions.cpp +++ b/tests/reductions.cpp @@ -96,3 +96,65 @@ TEST_F(Reductions, Min) { device->api->freePinnedMem(testResult); device->api->freeGlobMem(devVector); } + +/** + * The reductions above run over unsigned values only, where the neutral element of a maximum and + * the smallest representable value are the same thing. They are not for signed integers, and for + * floating point types numeric_limits::min() is the smallest positive normal value, so a + * maximum over negative data has to start below all of them to come out right. + */ +template +class SignedReductions : public BaseTestSuite { + protected: + void run(ReductionType type, const std::vector& host, T expected) { + auto* devVector = static_cast(device->api->allocGlobMem(sizeof(T) * host.size())); + device->api->copyTo(devVector, host.data(), sizeof(T) * host.size()); + + auto* result = static_cast(device->api->allocPinnedMem(sizeof(T))); + *result = T{0}; + + device->algorithms.reduceVector( + result, devVector, true, host.size(), type, device->api->getDefaultStream()); + device->api->syncDefaultStreamWithHost(); + + EXPECT_EQ(expected, *result); + + device->api->freePinnedMem(result); + device->api->freeGlobMem(devVector); + } +}; + +using SignedTypes = ::testing::Types; +TYPED_TEST_SUITE(SignedReductions, SignedTypes); + +TYPED_TEST(SignedReductions, maxOverNegativeValues) { + std::vector host(100000, TypeParam{-7}); + host[host.size() / 3] = TypeParam{-2}; + this->run(ReductionType::Max, host, TypeParam{-2}); +} + +TYPED_TEST(SignedReductions, minOverNegativeValues) { + std::vector host(100000, TypeParam{-7}); + host[host.size() / 3] = TypeParam{-11}; + this->run(ReductionType::Min, host, TypeParam{-11}); +} + +TYPED_TEST(SignedReductions, addOverNegativeValues) { + std::vector host(1000, TypeParam{-3}); + this->run(ReductionType::Add, host, TypeParam{-3000}); +} + +TEST_F(Reductions, emptyInput) { + auto* devVector = static_cast(device->api->allocGlobMem(sizeof(float))); + auto* result = static_cast(device->api->allocPinnedMem(sizeof(float))); + *result = 123.0F; + + // nothing to reduce, but the result is still initialized to the neutral element + device->algorithms.reduceVector( + result, devVector, true, 0, ReductionType::Add, device->api->getDefaultStream()); + device->api->syncDefaultStreamWithHost(); + EXPECT_EQ(0.0F, *result); + + device->api->freePinnedMem(result); + device->api->freeGlobMem(devVector); +} diff --git a/tests/streams.cpp b/tests/streams.cpp new file mode 100644 index 0000000..459c6c7 --- /dev/null +++ b/tests/streams.cpp @@ -0,0 +1,207 @@ +// SPDX-FileCopyrightText: 2026 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "BaseTestSuite.h" +#include "device.h" + +#include "gtest/gtest.h" +#include +#include +#include + +using namespace device; +using namespace ::testing; + +namespace { +constexpr std::size_t ArraySize = 1 << 14; + +// Enough queued work that the host is certain to run ahead of the producing stream. Without it, +// an ordering test passes whenever the producer happens to finish first, so a missing dependency +// shows up as an occasional failure rather than as a verdict. +constexpr int EnqueueDepth = 64; +} // namespace + +class Streams : public BaseTestSuite { + public: + void SetUp() override { + BaseTestSuite::SetUp(); + devArray = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); + streamA = device->api->createStream(); + streamB = device->api->createStream(); + } + + void TearDown() override { + device->api->destroyGenericStream(streamB); + device->api->destroyGenericStream(streamA); + device->api->freeGlobMem(devArray); + } + + protected: + std::vector download(void* stream) { + std::vector host(ArraySize, -1); + device->api->copyFromAsync(host.data(), devArray, ArraySize * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + return host; + } + + float* devArray{nullptr}; + void* streamA{nullptr}; + void* streamB{nullptr}; +}; + +TEST_F(Streams, anEventOrdersTwoStreams) { + auto* event = device->api->createEvent(); + + for (int i = 0; i < EnqueueDepth; ++i) { + device->algorithms.fillArray(devArray, 2.0F, ArraySize, streamA); + } + device->api->recordEventOnStream(event, streamA); + + device->api->syncStreamWithEvent(streamB, event); + device->algorithms.scaleArray(devArray, 7.0F, ArraySize, streamB); + + device->api->syncStreamWithHost(streamB); + + // a 7 here means the scale read the array before the fills wrote it + for (const auto value : download(streamB)) { + ASSERT_EQ(14.0F, value); + } + + device->api->destroyEvent(event); +} + +TEST_F(Streams, anEventCanBeRecordedAgain) { + // the stream runtime hands the same event out repeatedly, so re-recording one that has already + // been waited upon has to keep working + auto* event = device->api->createEvent(); + + for (int round = 1; round <= 4; ++round) { + for (int i = 0; i < EnqueueDepth; ++i) { + device->algorithms.fillArray(devArray, static_cast(round), ArraySize, streamA); + } + device->api->recordEventOnStream(event, streamA); + device->api->syncStreamWithEvent(streamB, event); + device->algorithms.scaleArray(devArray, 2.0F, ArraySize, streamB); + device->api->syncStreamWithHost(streamB); + + // twice the previous round's value means the scale overtook this round's fills + for (const auto value : download(streamB)) { + ASSERT_EQ(2.0F * static_cast(round), value) + << "in round " << round << " (twice the previous round's value would be " + << 4.0F * static_cast(round - 1) << ")"; + } + } + + device->api->destroyEvent(event); +} + +TEST_F(Streams, anEventIsCompleteOnceItsStreamIs) { + auto* event = device->api->createEvent(); + + device->algorithms.fillArray(devArray, 1.0F, ArraySize, streamA); + device->api->recordEventOnStream(event, streamA); + device->api->syncStreamWithHost(streamA); + + // only the direction after synchronizing is deterministic; whether the event is already + // complete beforehand depends on timing + EXPECT_TRUE(device->api->isEventCompleted(event)); + EXPECT_TRUE(device->api->isStreamWorkDone(streamA)); + + device->api->destroyEvent(event); +} + +TEST_F(Streams, aHostFunctionRunsInStreamOrder) { + std::vector staging(ArraySize, -1); + std::atomic sawFilledData{false}; + std::atomic ran{false}; + + device->algorithms.fillArray(devArray, 9.0F, ArraySize, streamA); + device->api->copyFromAsync(staging.data(), devArray, ArraySize * sizeof(float), streamA); + device->api->streamHostFunction(streamA, [&]() { + ran = true; + sawFilledData = (staging.front() == 9.0F) && (staging.back() == 9.0F); + }); + + device->api->syncStreamWithHost(streamA); + + EXPECT_TRUE(ran.load()); + // the callback must not observe the staging buffer before the copy that precedes it + EXPECT_TRUE(sawFilledData.load()); +} + +TEST_F(Streams, asyncAllocationsLiveOnTheStream) { + auto* scratch = + static_cast(device->api->allocMemAsync(ArraySize * sizeof(float), streamA)); + ASSERT_NE(nullptr, scratch); + + device->algorithms.fillArray(scratch, 4.0F, ArraySize, streamA); + + std::vector host(ArraySize, -1); + device->api->copyFromAsync(host.data(), scratch, ArraySize * sizeof(float), streamA); + + // the free is ordered behind the copy on the same stream + device->api->freeMemAsync(scratch, streamA); + device->api->syncStreamWithHost(streamA); + + for (const auto value : host) { + ASSERT_EQ(4.0F, value); + } +} + +TEST_F(Streams, aDestroyedStreamIsForgotten) { + // A device-wide synchronization walks every stream the backend knows about. A stream that was + // destroyed therefore has to be off that list, or the walk runs into freed memory - long after + // the code that destroyed it, which is what makes this kind of fault hard to place. + auto* scratch = device->api->createStream(); + device->algorithms.fillArray(devArray, 1.0F, ArraySize, scratch); + device->api->syncStreamWithHost(scratch); + device->api->destroyGenericStream(scratch); + + device->api->syncDevice(); + + for (const auto value : download(streamA)) { + ASSERT_EQ(1.0F, value); + } +} + +TEST_F(Streams, workOnSeparateStreamsStaysSeparate) { + auto* other = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); + + device->algorithms.fillArray(devArray, 1.0F, ArraySize, streamA); + device->algorithms.fillArray(other, 2.0F, ArraySize, streamB); + + device->api->syncStreamWithHost(streamA); + device->api->syncStreamWithHost(streamB); + + std::vector hostOther(ArraySize, -1); + device->api->copyFromAsync(hostOther.data(), other, ArraySize * sizeof(float), streamB); + device->api->syncStreamWithHost(streamB); + + for (const auto value : download(streamA)) { + ASSERT_EQ(1.0F, value); + } + for (const auto value : hostOther) { + ASSERT_EQ(2.0F, value); + } + + device->api->freeGlobMem(other); +} + +TEST_F(Streams, streamsCanBeGivenAPriority) { + // 0 is the lowest priority the device offers, 1 the highest, and the default is whatever the + // runtime picks; all three have to give a stream that works + for (const double priority : {0.0, 0.5, 1.0}) { + auto* stream = device->api->createStream(priority); + ASSERT_NE(nullptr, stream) << "at priority " << priority; + + device->algorithms.fillArray(devArray, 6.0F, ArraySize, stream); + device->api->syncStreamWithHost(stream); + + for (const auto value : download(stream)) { + ASSERT_EQ(6.0F, value) << "at priority " << priority; + } + + device->api->destroyGenericStream(stream); + } +}