From 2a617e05a0a4072c157085f311aa6d9e8588e4e8 Mon Sep 17 00:00:00 2001 From: Elliott Slaughter Date: Fri, 28 Aug 2026 16:11:53 -0700 Subject: [PATCH] Implement Realm reductions for CUDA. --- .flake/pkgs/realm.nix | 6 +- CMakeLists.txt | 8 +- cmake/cuda.cmake | 1 + cmake/flexflow-utils.cmake | 18 +- lib/realm-execution/CMakeLists.txt | 4 + .../redops/realm_redop_registry.cc | 543 +---------- .../redops/realm_redop_registry.cu | 885 ++++++++++++++++++ .../src/realm-execution/redops/redop_id_t.cc | 2 +- 8 files changed, 919 insertions(+), 548 deletions(-) create mode 100644 lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cu diff --git a/.flake/pkgs/realm.nix b/.flake/pkgs/realm.nix index f66da107df..6e5708ad32 100644 --- a/.flake/pkgs/realm.nix +++ b/.flake/pkgs/realm.nix @@ -13,13 +13,13 @@ in stdenv.mkDerivation rec { pname = "realm"; - version = "2026-07-01"; + version = "2026-08-28"; src = fetchFromGitHub { owner = "StanfordLegion"; repo = "realm"; - rev = "dba83452d4a211e3a0c0e5d4849c26ce7b56bba3"; - sha256 = "sha256-dLGmUgupDX0dDErZbuP5ipj2Cwf+rOhVGC/+angRyBY="; + rev = "15b3e9b68b65d698d237766d37bc36522b8ecf0b"; + sha256 = "sha256-HXNRKzOIUok8FAuhJF1sOW0+Hv7rRfRDoQNt+wAvOnc="; }; nativeBuildInputs = [ diff --git a/CMakeLists.txt b/CMakeLists.txt index 4d62faf688..0dffa77847 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -51,15 +51,17 @@ option(FF_BUILD_ARG_PARSER "build command line argument parser" OFF) option(FF_BUILD_BIN_EXPORT_MODEL_ARCH "build export-model-arch utility" ON) option(FF_BUILD_BIN_RUN_MODEL "build run-model binary" ON) -include(cuda) -include(cudnn) -include(nccl) if (FF_USE_CODE_COVERAGE) include(CodeCoverage) append_coverage_compiler_flags() endif() # set_property(CACHE FF_GPU_BACKEND PROPERTY STRINGS ${FF_GPU_BACKENDS}) +# Load after CodeCoverage because CodeCoverage blows up on NVCC compiler +include(cuda) +include(cudnn) +include(nccl) + include(json) include(expected) include(spdlog) diff --git a/cmake/cuda.cmake b/cmake/cuda.cmake index e12ac52589..9bc4109f17 100644 --- a/cmake/cuda.cmake +++ b/cmake/cuda.cmake @@ -1,5 +1,6 @@ include(aliasing) +enable_language(CUDA) find_package(CUDAToolkit 11.7 REQUIRED) alias_library(deps::cublas CUDA::cublas) diff --git a/cmake/flexflow-utils.cmake b/cmake/flexflow-utils.cmake index 5dccd5272d..a7da282543 100644 --- a/cmake/flexflow-utils.cmake +++ b/cmake/flexflow-utils.cmake @@ -56,6 +56,14 @@ function(ff_set_cxx_properties target) ) endfunction() +function(ff_set_cuda_properties target) + set_target_properties(${target} + PROPERTIES + CUDA_STANDARD 17 + CUDA_STANDARD_REQUIRED YES + ) +endfunction() + function(ff_add_library) ff_parse_args( PREFIX @@ -63,6 +71,7 @@ function(ff_add_library) ARGS NAME VARIADIC_ARGS + LANGUAGES SRC_PATTERNS PUBLIC_INCLUDE PRIVATE_INCLUDE @@ -72,7 +81,11 @@ function(ff_add_library) ${ARGN} ) - project(${FF_LIBRARY_NAME}) + if(NOT FF_LIBRARY_LANGUAGES) + project(${FF_LIBRARY_NAME}) + else() + project(${FF_LIBRARY_NAME} LANGUAGES ${FF_LIBRARY_LANGUAGES}) + endif() file(GLOB_RECURSE SRC CONFIGURE_DEPENDS LIST_DIRECTORIES False @@ -99,6 +112,9 @@ function(ff_add_library) ) define_ff_vars(${FF_LIBRARY_NAME}) ff_set_cxx_properties(${FF_LIBRARY_NAME}) + if(CUDA IN_LIST FF_LIBRARY_LANGUAGES) + ff_set_cuda_properties(${FF_LIBRARY_NAME}) + endif() endfunction() function(ff_add_test_executable) diff --git a/lib/realm-execution/CMakeLists.txt b/lib/realm-execution/CMakeLists.txt index 49fbcfa4e0..2a8f3206e8 100644 --- a/lib/realm-execution/CMakeLists.txt +++ b/lib/realm-execution/CMakeLists.txt @@ -1,8 +1,12 @@ ff_add_library( NAME realm-execution + LANGUAGES + CXX + CUDA SRC_PATTERNS src/*.cc + src/*.cu PUBLIC_INCLUDE include/ PRIVATE_INCLUDE diff --git a/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cc b/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cc index ab3304836a..3b4350bc11 100644 --- a/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cc +++ b/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cc @@ -1,540 +1,3 @@ -#include "realm-execution/redops/realm_redop_registry.h" -#include "realm-execution/redops/redop_id_t.h" - -namespace FlexFlow { - -// Reduction operators and related infrastructure borrowed from Legion. We -// maintain the Legion naming scheme to maximizing compatibility with the -// existing code, despite not otherwise relying or using Legion in any way. -// https://gitlab.com/StanfordLegion/legion/-/blob/5263aeff477fb94239c50d9306d58c4244e9fc38/runtime/legion/api/redop.inl#L31 -#if !defined(__cpp_lib_atomic_ref) || (__cpp_lib_atomic_ref < 201806L) -// We only need this crap if we're using a version of c++ < 20 -// Starting with c++20 we can do all this the right way with atomic_ref -namespace TypePunning { -// The tenth circle of hell is reserved for members of the C++ committee -// that decided to deviate from C's support for type punning unions. -// Add on to it the fact that it took them 9 fucking years to realize -// that they needed std::atomic_ref and it's plain to see they are all -// just a bunch of idiots that should never be allowed near a programming -// language standard ever again. They've clearly never written lock-free -// code in their lives. -template -class Pointer { -public: - Pointer(void *p) : pointer(convert(p)) {} - static inline T *convert(void *p) { - T *ptr = nullptr; - static_assert(sizeof(ptr) == sizeof(p)); - memcpy(&ptr, &p, sizeof(p)); - return ptr; - } - inline operator T *(void) const { - return (T *)pointer; - } - inline T operator*(void) const { - return *pointer; - } - inline T operator[](size_t off) const { - return pointer[off]; - } - -private: - T volatile *const pointer; -}; -template -class AlignedPointer { -public: - AlignedPointer(void *p) : off(align(p)), pointer(convert(p, off)) {} - static inline T *convert(void *p, size_t off) { - uint8_t *p1 = nullptr; - static_assert(sizeof(p1) == sizeof(p)); - memcpy(&p1, &p, sizeof(p)); - p1 = p1 - off; - T *p2 = nullptr; - static_assert(sizeof(p1) == sizeof(p2)); - memcpy(&p2, &p1, sizeof(p1)); - return p2; - } - static inline size_t align(void *p) { - uintptr_t ptr; - static_assert(sizeof(ptr) == sizeof(p)); - memcpy(&ptr, &p, sizeof(ptr)); - return ptr % ALIGNMENT; - } - inline operator T *(void) const { - return (T *)pointer; - } - inline T operator*(void) const { - return *pointer; - } - inline size_t offset(void) const { - return off; - } - -private: - size_t off; - T volatile *const pointer; -}; -template -class Alias { -public: - inline void load(Pointer const &pointer, size_t off = 0) { - T1 value = pointer[off]; - memcpy(buffer, (void *)&value, sizeof(T1)); - } - template - inline void load(AlignedPointer const &pointer) { - T1 value = *pointer; - memcpy(buffer, (void *)&value, sizeof(T1)); - } - inline T1 as_one(void) const { - T1 result; - memcpy((void *)&result, buffer, sizeof(result)); - return result; - } - inline T2 as_two(void) const { - T2 result; - memcpy((void *)&result, buffer, sizeof(result)); - return result; - } - inline Alias &operator=(T2 rhs) { - memcpy(buffer, (void *)&rhs, sizeof(rhs)); - return *this; - } - -private: - // Make this one private so it is can never be called - inline Alias &operator=(T1 rhs) { - memcpy(buffer, (void *)&rhs, sizeof(rhs)); - return *this; - } - static_assert(sizeof(T1) == sizeof(T2)); - uint8_t buffer[sizeof(T1)]; -}; -}; // namespace TypePunning -#endif - -// Define a prefix for annotating functions for CUDA compilation -#if defined(__CUDACC__) || defined(__HIPCC__) -#define __LEGION_CUDA_HD__ __host__ __device__ -#else -#define __LEGION_CUDA_HD__ -#endif - -template -class SumReduction { - // Empty definition - // Specializations provided for each type -}; - -template <> -class SumReduction { -public: - typedef bool LHS; - typedef bool RHS; - - static constexpr bool identity = false; - - template - __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); - template - __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); -}; - -template <> -class SumReduction { -public: - typedef int32_t LHS; - typedef int32_t RHS; - - static constexpr int32_t identity = 0; - - template - __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); - template - __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); -}; - -template <> -class SumReduction { -public: - typedef int64_t LHS; - typedef int64_t RHS; - - static constexpr int64_t identity = 0; - - template - __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); - template - __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); -}; - -template <> -class SumReduction { -public: - typedef float LHS; - typedef float RHS; - - static constexpr float identity = 0.f; - - template - __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); - template - __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); -}; - -template <> -class SumReduction { -public: - typedef double LHS; - typedef double RHS; - - static constexpr double identity = 0.0; - - template - __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); - template - __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); -}; - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { - lhs = lhs || rhs; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - // GPU atomics need 4 byte alignment - const uintptr_t unaligned = reinterpret_cast(&lhs); - unsigned const offset = unaligned % sizeof(unsigned int); - const uintptr_t aligned = unaligned - offset; - unsigned int *ptr = reinterpret_cast(aligned); - unsigned int newval = *ptr, oldval; - do { - RHS previous = __uint2bool(newval, offset); - RHS next = previous || rhs; - oldval = newval; - newval = __bool2uint(newval, next, offset); - newval = atomicCAS(ptr, oldval, newval); - } while (oldval != newval); -#else -#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) - std::atomic_ref atomic(lhs); - RHS oldval = atomic.load(); - RHS newval; - do { - newval = oldval || rhs; - } while (!atomic.compare_exchange_weak(oldval, newval)); -#else - // No atomic logical operations so use compare and swap - TypePunning::Alias oldval, newval; - TypePunning::Pointer pointer((void *)&lhs); - do { - oldval.load(pointer); - newval = oldval.as_two() || rhs; - } while (!__sync_bool_compare_and_swap( - (int8_t *)pointer, oldval.as_one(), newval.as_one())); -#endif -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { - rhs1 = rhs1 || rhs2; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - // GPU atomics need 4 byte alignment - const uintptr_t unaligned = reinterpret_cast(&rhs1); - unsigned const offset = unaligned % sizeof(unsigned int); - const uintptr_t aligned = unaligned - offset; - unsigned int *ptr = reinterpret_cast(aligned); - unsigned int newval = *ptr, oldval; - do { - RHS previous = __uint2bool(newval, offset); - RHS next = previous || rhs2; - oldval = newval; - newval = __bool2uint(newval, next, offset); - newval = atomicCAS(ptr, oldval, newval); - } while (oldval != newval); -#else -#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) - std::atomic_ref atomic(rhs1); - RHS oldval = atomic.load(); - RHS newval; - do { - newval = oldval || rhs2; - } while (!atomic.compare_exchange_weak(oldval, newval)); -#else - // No atomic logical operations so use compare and swap - TypePunning::Alias oldval, newval; - TypePunning::Pointer pointer((void *)&rhs1); - do { - oldval.load(pointer); - newval = oldval.as_two() || rhs2; - } while (!__sync_bool_compare_and_swap( - (int8_t *)pointer, oldval.as_one(), newval.as_one())); -#endif -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { - lhs += rhs; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - atomicAdd(&lhs, rhs); -#else - __sync_fetch_and_add(&lhs, rhs); -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { - rhs1 += rhs2; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - atomicAdd(&rhs1, rhs2); -#else - __sync_fetch_and_add(&rhs1, rhs2); -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { - lhs += rhs; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - // Apparently there is no signed 64bit int atomic yet - RHS newval = lhs, oldval; - // Type punning like this is illegal in C++ but the - // CUDA manual has an example just like it so fuck it - unsigned long long int *ptr = (unsigned long long int *)&lhs; - do { - oldval = newval; - newval += rhs; - newval = __ulonglong_as_longlong(atomicCAS( - ptr, __longlong_as_ulonglong(oldval), __longlong_as_ulonglong(newval))); - } while (oldval != newval); -#else - __sync_fetch_and_add(&lhs, rhs); -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { - rhs1 += rhs2; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - // Apparently there is no signed 64bit int atomic yet - RHS newval = rhs1, oldval; - // Type punning like this is illegal in C++ but the - // CUDA manual has an example just like it so fuck it - unsigned long long int *ptr = (unsigned long long int *)&rhs1; - do { - oldval = newval; - newval += rhs2; - newval = __ulonglong_as_longlong(atomicCAS( - ptr, __longlong_as_ulonglong(oldval), __longlong_as_ulonglong(newval))); - } while (oldval != newval); -#else - __sync_fetch_and_add(&rhs1, rhs2); -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { - lhs += rhs; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - atomicAdd(&lhs, rhs); -#else -#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) - std::atomic_ref atomic(lhs); - RHS oldval = atomic.load(); - RHS newval; - do { - newval = oldval + rhs; - } while (!atomic.compare_exchange_weak(oldval, newval)); -#else - // No atomic floating point operations so use compare and swap - TypePunning::Alias oldval, newval; - TypePunning::Pointer pointer((void *)&lhs); - do { - oldval.load(pointer); - newval = oldval.as_two() + rhs; - } while (!__sync_bool_compare_and_swap( - (int32_t *)pointer, oldval.as_one(), newval.as_one())); -#endif -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { - rhs1 += rhs2; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - atomicAdd(&rhs1, rhs2); -#else -#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) - std::atomic_ref atomic(rhs1); - RHS oldval = atomic.load(); - RHS newval; - do { - newval = oldval + rhs2; - } while (!atomic.compare_exchange_weak(oldval, newval)); -#else - // No atomic floating point operations so use compare and swap - TypePunning::Alias oldval, newval; - TypePunning::Pointer pointer((void *)&rhs1); - do { - oldval.load(pointer); - newval = oldval.as_two() + rhs2; - } while (!__sync_bool_compare_and_swap( - (int32_t *)pointer, oldval.as_one(), newval.as_one())); -#endif -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { - lhs += rhs; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, - RHS rhs) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) -#if (__CUDA_ARCH__ >= 600) || defined(__HIP_DEVICE_COMPILE__) - atomicAdd(&lhs, rhs); -#else - RHS newval = lhs, oldval; - // Type punning like this is illegal in C++ but the - // CUDA manual has an example just like it so fuck it - unsigned long long int *ptr = (unsigned long long int *)&lhs; - do { - oldval = newval; - newval += rhs; - newval = __ulonglong_as_double(atomicCAS( - ptr, __double_as_ulonglong(oldval), __double_as_ulonglong(newval))); - } while (oldval != newval); -#endif -#else -#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) - std::atomic_ref atomic(lhs); - RHS oldval = atomic.load(); - RHS newval; - do { - newval = oldval + rhs; - } while (!atomic.compare_exchange_weak(oldval, newval)); -#else - // No atomic floating point operations so use compare and swap - TypePunning::Alias oldval, newval; - TypePunning::Pointer pointer((void *)&lhs); - do { - oldval.load(pointer); - newval = oldval.as_two() + rhs; - } while (!__sync_bool_compare_and_swap( - (int64_t *)pointer, oldval.as_one(), newval.as_one())); -#endif -#endif -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { - rhs1 += rhs2; -} - -template <> -__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, - RHS rhs2) { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) -#if (__CUDA_ARCH__ >= 600) || defined(__HIP_DEVICE_COMPILE__) - atomicAdd(&rhs1, rhs2); -#else - RHS newval = rhs1, oldval; - // Type punning like this is illegal in C++ but the - // CUDA manual has an example just like it so fuck it - unsigned long long int *ptr = (unsigned long long int *)&rhs1; - do { - oldval = newval; - newval += rhs2; - newval = __ulonglong_as_double(atomicCAS( - ptr, __double_as_ulonglong(oldval), __double_as_ulonglong(newval))); - } while (oldval != newval); -#endif -#else -#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) - std::atomic_ref atomic(rhs1); - RHS oldval = atomic.load(); - RHS newval; - do { - newval = oldval + rhs2; - } while (!atomic.compare_exchange_weak(oldval, newval)); -#else - // No atomic floating point operations so use compare and swap - TypePunning::Alias oldval, newval; - TypePunning::Pointer pointer((void *)&rhs1); - do { - oldval.load(pointer); - newval = oldval.as_two() + rhs2; - } while (!__sync_bool_compare_and_swap( - (int64_t *)pointer, oldval.as_one(), newval.as_one())); -#endif -#endif -} - -void register_all_redops(Realm::Runtime rt) { - // Registration is synchronous, so no need to capture events here - rt.register_reduction>( - get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_BOOL_REDOP_ID)); - rt.register_reduction>( - get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_INT32_REDOP_ID)); - rt.register_reduction>( - get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_INT64_REDOP_ID)); - rt.register_reduction>( - get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_FLOAT_REDOP_ID)); - rt.register_reduction>( - get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_DOUBLE_REDOP_ID)); -} - -} // namespace FlexFlow +// See implementation in realm_redop_registry.cu +// +// This file is here to make proj layout checks happy diff --git a/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cu b/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cu new file mode 100644 index 0000000000..4614c71e40 --- /dev/null +++ b/lib/realm-execution/src/realm-execution/redops/realm_redop_registry.cu @@ -0,0 +1,885 @@ +#include "realm-execution/redops/realm_redop_registry.h" +#include "realm-execution/redops/redop_id_t.h" + +namespace FlexFlow { + +// Reduction operators and related infrastructure borrowed from Legion. We +// maintain the Legion naming scheme to maximizing compatibility with the +// existing code, despite not otherwise relying or using Legion in any way. +// https://gitlab.com/StanfordLegion/legion/-/blob/5263aeff477fb94239c50d9306d58c4244e9fc38/runtime/legion/api/redop.inl#L31 +#if !defined(__cpp_lib_atomic_ref) || (__cpp_lib_atomic_ref < 201806L) +// We only need this crap if we're using a version of c++ < 20 +// Starting with c++20 we can do all this the right way with atomic_ref +namespace TypePunning { +// The tenth circle of hell is reserved for members of the C++ committee +// that decided to deviate from C's support for type punning unions. +// Add on to it the fact that it took them 9 fucking years to realize +// that they needed std::atomic_ref and it's plain to see they are all +// just a bunch of idiots that should never be allowed near a programming +// language standard ever again. They've clearly never written lock-free +// code in their lives. +template +class Pointer { +public: + Pointer(void *p) : pointer(convert(p)) {} + static inline T *convert(void *p) { + T *ptr = nullptr; + static_assert(sizeof(ptr) == sizeof(p)); + memcpy(&ptr, &p, sizeof(p)); + return ptr; + } + inline operator T *(void) const { + return (T *)pointer; + } + inline T operator*(void) const { + return *pointer; + } + inline T operator[](size_t off) const { + return pointer[off]; + } + +private: + T volatile *const pointer; +}; +template +class AlignedPointer { +public: + AlignedPointer(void *p) : off(align(p)), pointer(convert(p, off)) {} + static inline T *convert(void *p, size_t off) { + uint8_t *p1 = nullptr; + static_assert(sizeof(p1) == sizeof(p)); + memcpy(&p1, &p, sizeof(p)); + p1 = p1 - off; + T *p2 = nullptr; + static_assert(sizeof(p1) == sizeof(p2)); + memcpy(&p2, &p1, sizeof(p1)); + return p2; + } + static inline size_t align(void *p) { + uintptr_t ptr; + static_assert(sizeof(ptr) == sizeof(p)); + memcpy(&ptr, &p, sizeof(ptr)); + return ptr % ALIGNMENT; + } + inline operator T *(void) const { + return (T *)pointer; + } + inline T operator*(void) const { + return *pointer; + } + inline size_t offset(void) const { + return off; + } + +private: + size_t off; + T volatile *const pointer; +}; +template +class Alias { +public: + inline void load(Pointer const &pointer, size_t off = 0) { + T1 value = pointer[off]; + memcpy(buffer, (void *)&value, sizeof(T1)); + } + template + inline void load(AlignedPointer const &pointer) { + T1 value = *pointer; + memcpy(buffer, (void *)&value, sizeof(T1)); + } + inline T1 as_one(void) const { + T1 result; + memcpy((void *)&result, buffer, sizeof(result)); + return result; + } + inline T2 as_two(void) const { + T2 result; + memcpy((void *)&result, buffer, sizeof(result)); + return result; + } + inline Alias &operator=(T2 rhs) { + memcpy(buffer, (void *)&rhs, sizeof(rhs)); + return *this; + } + +private: + // Make this one private so it is can never be called + inline Alias &operator=(T1 rhs) { + memcpy(buffer, (void *)&rhs, sizeof(rhs)); + return *this; + } + static_assert(sizeof(T1) == sizeof(T2)); + uint8_t buffer[sizeof(T1)]; +}; +}; // namespace TypePunning +#endif + +// Define a prefix for annotating functions for CUDA compilation +#if defined(__CUDACC__) || defined(__HIPCC__) +#define __LEGION_CUDA_HD__ __host__ __device__ +#else +#define __LEGION_CUDA_HD__ +#endif + +#if defined(__CUDACC__) || defined(__HIPCC__) + // We have these functions here because calling memcpy (per the +// insistence of the idiots on the C++ standards committee) on the +// GPU is a terrible idea since it will spill the data out of registers +// and into local memory in order to do the memcpy. Hence we tell +// the compiler exactly what we mean using bit operations and inline PTX. + +__device__ __forceinline__ bool __uint2bool(unsigned int value, + unsigned offset) { + value = value >> (8 * offset); + return ((value & 0xFF) != 0); +} + +__device__ __forceinline__ unsigned int + __bool2uint(unsigned int previous, bool value, unsigned offset) { + unsigned int next = value; + next = next << (8 * offset); + unsigned int mask = 0xFF; + mask = mask << (8 * offset); + previous = previous & (~mask); + return previous | next; +} + +__device__ __forceinline__ uint8_t __uint2ubyte(unsigned int value, + unsigned offset) { + value = value >> (8 * offset); + return uint8_t(value & 0xFF); +} + +__device__ __forceinline__ unsigned int + __ubyte2uint(unsigned int previous, uint8_t value, unsigned offset) { + unsigned int next = value; + next = next << (8 * offset); + unsigned int mask = 0xFF; + mask = mask << (8 * offset); + previous = previous & (~mask); + return previous | next; +} + +__device__ __forceinline__ int8_t __int2byte(int value, unsigned offset) { + value = value >> (8 * offset); + return int8_t(value & 0xFF); +} + +__device__ __forceinline__ int + __byte2int(int previous, int8_t value, unsigned offset) { + int next = value; + next = next << (8 * offset); + unsigned int mask = 0xFF; + mask = mask << (8 * offset); + previous = previous & (~mask); + return previous | next; +} + +__device__ __forceinline__ unsigned short int + __short_as_ushort(short int value) { +#ifdef __HIPCC__ + union { + short int as_signed; + unsigned short int as_unsigned; + } val; + val.as_signed = value; + return val.as_unsigned; +#else + unsigned short int result; + asm("mov.b16 %0, %1;" : "=h"(result) : "h"(value)); + return result; +#endif +} + +__device__ __forceinline__ short int + __ushort_as_short(unsigned short int value) { +#ifdef __HIPCC__ + union { + short int as_signed; + unsigned short int as_unsigned; + } val; + val.as_unsigned = value; + return val.as_signed; +#else + short int result; + asm("mov.b16 %0, %1;" : "=h"(result) : "h"(value)); + return result; +#endif +} + +__device__ __forceinline__ unsigned int + __hiloushort2uint(unsigned short int hi, unsigned short int lo) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + ushort2 as_short; + } val; + val.as_short.x = lo; + val.as_short.y = hi; + return val.as_int; +#else + unsigned int result; + asm("mov.b32 %0, {%1,%2};" : "=r"(result) : "h"(lo), "h"(hi)); + return result; +#endif +} + +__device__ __forceinline__ unsigned short int + __uint2loushort(unsigned int value) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + ushort2 as_short; + } val; + val.as_int = value; + return val.as_short.x; +#else + unsigned short int lo, hi; + asm("mov.b32 {%0,%1}, %2;" : "=h"(lo), "=h"(hi) : "r"(value)); + return lo + 0 * hi; +#endif +} + +__device__ __forceinline__ unsigned short int + __uint2hiushort(unsigned int value) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + ushort2 as_short; + } val; + val.as_int = value; + return val.as_short.y; +#else + unsigned short int lo, hi; + asm("mov.b32 {%0,%1}, %2;" : "=h"(lo), "=h"(hi) : "r"(value)); + return hi + 0 * lo; +#endif +} + +__device__ __forceinline__ unsigned int __hiloshort2uint(short int hi, + short int lo) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + short2 as_short; + } val; + val.as_short.x = lo; + val.as_short.y = hi; + return val.as_int; +#else + unsigned int result; + asm("mov.b32 %0, {%1,%2};" : "=r"(result) : "h"(lo), "h"(hi)); + return result; +#endif +} + +__device__ __forceinline__ short int __uint2loshort(unsigned int value) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + short2 as_short; + } val; + val.as_int = value; + return val.as_short.x; +#else + short int lo, hi; + asm("mov.b32 {%0,%1}, %2;" : "=h"(lo), "=h"(hi) : "r"(value)); + return lo + 0 * hi; +#endif +} + +__device__ __forceinline__ short int __uint2hishort(unsigned int value) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + short2 as_short; + } val; + val.as_int = value; + return val.as_short.y; +#else + short int lo, hi; + asm("mov.b32 {%0,%1}, %2;" : "=h"(lo), "=h"(hi) : "r"(value)); + return hi + 0 * lo; +#endif +} + +#ifdef LEGION_REDOP_HALF +__device__ __forceinline__ unsigned int __hilohalf2uint(__half hi, __half lo) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + short2 as_short; + } val; + val.as_short.x = __half_as_short(lo); + val.as_short.y = __half_as_short(hi); + return val.as_int; +#else + unsigned int result; + asm("mov.b32 %0, {%1,%2};" + : "=r"(result) + : "h"(__half_as_short(lo)), "h"(__half_as_short(hi))); + return result; +#endif +} + +__device__ __forceinline__ __half __uint2hihalf(unsigned int value) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + short2 as_short; + } val; + val.as_int = value; + return __short_as_half(val.as_short.y); +#else + short int lo, hi; + asm("mov.b32 {%0,%1}, %2;" : "=h"(lo), "=h"(hi) : "r"(value)); + return __short_as_half(hi) + __half(0) * __short_as_half(lo); +#endif +} + +__device__ __forceinline__ __half __uint2lohalf(unsigned int value) { +#ifdef __HIPCC__ + union { + unsigned int as_int; + short2 as_short; + } val; + val.as_int = value; + return __short_as_half(val.as_short.x); +#else + short int lo, hi; + asm("mov.b32 {%0,%1}, %2;" : "=h"(lo), "=h"(hi) : "r"(value)); + return __short_as_half(lo) + __half(0) * __short_as_half(hi); +#endif +} +#endif + +__device__ __forceinline__ unsigned int __int_as_uint(int value) { +#ifdef __HIPCC__ + union { + int as_signed; + unsigned int as_unsigned; + } val; + val.as_signed = value; + return val.as_unsigned; +#else + unsigned int result; + asm("mov.b32 %0, %1;" : "=r"(result) : "r"(value)); + return result; +#endif +} + +__device__ __forceinline__ int __uint_as_int(unsigned int value) { +#ifdef __HIPCC__ + union { + int as_signed; + unsigned int as_unsigned; + } val; + val.as_unsigned = value; + return val.as_signed; +#else + int result; + asm("mov.b32 %0, %1;" : "=r"(result) : "r"(value)); + return result; +#endif +} + +__device__ __forceinline__ unsigned long long + __longlong_as_ulonglong(long long value) { +#ifdef __HIPCC__ + union { + long long as_signed; + unsigned long long as_unsigned; + } val; + val.as_signed = value; + return val.as_unsigned; +#else + unsigned long long result; + asm("mov.b64 %0, %1;" : "=l"(result) : "l"(value)); + return result; +#endif +} + +__device__ __forceinline__ long long + __ulonglong_as_longlong(unsigned long long value) { +#ifdef __HIPCC__ + union { + long long as_signed; + unsigned long long as_unsigned; + } val; + val.as_unsigned = value; + return val.as_signed; +#else + long long result; + asm("mov.b64 %0, %1;" : "=l"(result) : "l"(value)); + return result; +#endif +} + +__device__ __forceinline__ double + __ulonglong_as_double(unsigned long long value) { +#ifdef __HIPCC__ + union { + unsigned long long as_int; + double as_float; + } val; + val.as_int = value; + return val.as_float; +#else + double result; + asm("mov.b64 %0, %1;" : "=d"(result) : "l"(value)); + return result; +#endif +} + +__device__ __forceinline__ unsigned long long + __double_as_ulonglong(double value) { +#ifdef __HIPCC__ + union { + unsigned long long as_int; + double as_float; + } val; + val.as_float = value; + return val.as_int; +#else + unsigned long long result; + asm("mov.b64 %0, %1;" : "=l"(result) : "d"(value)); + return result; +#endif +} +#endif + +template +class SumReduction { + // Empty definition + // Specializations provided for each type +}; + +template <> +class SumReduction { +public: + typedef bool LHS; + typedef bool RHS; + + static constexpr bool identity = false; + + template + __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); + template + __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); +}; + +template <> +class SumReduction { +public: + typedef int32_t LHS; + typedef int32_t RHS; + + static constexpr int32_t identity = 0; + + template + __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); + template + __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); +}; + +template <> +class SumReduction { +public: + typedef int64_t LHS; + typedef int64_t RHS; + + static constexpr int64_t identity = 0; + + template + __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); + template + __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); +}; + +template <> +class SumReduction { +public: + typedef float LHS; + typedef float RHS; + + static constexpr float identity = 0.f; + + template + __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); + template + __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); +}; + +template <> +class SumReduction { +public: + typedef double LHS; + typedef double RHS; + + static constexpr double identity = 0.0; + + template + __LEGION_CUDA_HD__ static void apply(LHS &lhs, RHS rhs); + template + __LEGION_CUDA_HD__ static void fold(RHS &rhs1, RHS rhs2); +}; + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { + lhs = lhs || rhs; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + // GPU atomics need 4 byte alignment + const uintptr_t unaligned = reinterpret_cast(&lhs); + unsigned const offset = unaligned % sizeof(unsigned int); + const uintptr_t aligned = unaligned - offset; + unsigned int *ptr = reinterpret_cast(aligned); + unsigned int newval = *ptr, oldval; + do { + RHS previous = __uint2bool(newval, offset); + RHS next = previous || rhs; + oldval = newval; + newval = __bool2uint(newval, next, offset); + newval = atomicCAS(ptr, oldval, newval); + } while (oldval != newval); +#else +#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) + std::atomic_ref atomic(lhs); + RHS oldval = atomic.load(); + RHS newval; + do { + newval = oldval || rhs; + } while (!atomic.compare_exchange_weak(oldval, newval)); +#else + // No atomic logical operations so use compare and swap + TypePunning::Alias oldval, newval; + TypePunning::Pointer pointer((void *)&lhs); + do { + oldval.load(pointer); + newval = oldval.as_two() || rhs; + } while (!__sync_bool_compare_and_swap( + (int8_t *)pointer, oldval.as_one(), newval.as_one())); +#endif +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { + rhs1 = rhs1 || rhs2; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + // GPU atomics need 4 byte alignment + const uintptr_t unaligned = reinterpret_cast(&rhs1); + unsigned const offset = unaligned % sizeof(unsigned int); + const uintptr_t aligned = unaligned - offset; + unsigned int *ptr = reinterpret_cast(aligned); + unsigned int newval = *ptr, oldval; + do { + RHS previous = __uint2bool(newval, offset); + RHS next = previous || rhs2; + oldval = newval; + newval = __bool2uint(newval, next, offset); + newval = atomicCAS(ptr, oldval, newval); + } while (oldval != newval); +#else +#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) + std::atomic_ref atomic(rhs1); + RHS oldval = atomic.load(); + RHS newval; + do { + newval = oldval || rhs2; + } while (!atomic.compare_exchange_weak(oldval, newval)); +#else + // No atomic logical operations so use compare and swap + TypePunning::Alias oldval, newval; + TypePunning::Pointer pointer((void *)&rhs1); + do { + oldval.load(pointer); + newval = oldval.as_two() || rhs2; + } while (!__sync_bool_compare_and_swap( + (int8_t *)pointer, oldval.as_one(), newval.as_one())); +#endif +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { + lhs += rhs; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + atomicAdd(&lhs, rhs); +#else + __sync_fetch_and_add(&lhs, rhs); +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { + rhs1 += rhs2; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + atomicAdd(&rhs1, rhs2); +#else + __sync_fetch_and_add(&rhs1, rhs2); +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { + lhs += rhs; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + // Apparently there is no signed 64bit int atomic yet + RHS newval = lhs, oldval; + // Type punning like this is illegal in C++ but the + // CUDA manual has an example just like it so fuck it + unsigned long long int *ptr = (unsigned long long int *)&lhs; + do { + oldval = newval; + newval += rhs; + newval = __ulonglong_as_longlong(atomicCAS( + ptr, __longlong_as_ulonglong(oldval), __longlong_as_ulonglong(newval))); + } while (oldval != newval); +#else + __sync_fetch_and_add(&lhs, rhs); +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { + rhs1 += rhs2; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + // Apparently there is no signed 64bit int atomic yet + RHS newval = rhs1, oldval; + // Type punning like this is illegal in C++ but the + // CUDA manual has an example just like it so fuck it + unsigned long long int *ptr = (unsigned long long int *)&rhs1; + do { + oldval = newval; + newval += rhs2; + newval = __ulonglong_as_longlong(atomicCAS( + ptr, __longlong_as_ulonglong(oldval), __longlong_as_ulonglong(newval))); + } while (oldval != newval); +#else + __sync_fetch_and_add(&rhs1, rhs2); +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { + lhs += rhs; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + atomicAdd(&lhs, rhs); +#else +#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) + std::atomic_ref atomic(lhs); + RHS oldval = atomic.load(); + RHS newval; + do { + newval = oldval + rhs; + } while (!atomic.compare_exchange_weak(oldval, newval)); +#else + // No atomic floating point operations so use compare and swap + TypePunning::Alias oldval, newval; + TypePunning::Pointer pointer((void *)&lhs); + do { + oldval.load(pointer); + newval = oldval.as_two() + rhs; + } while (!__sync_bool_compare_and_swap( + (int32_t *)pointer, oldval.as_one(), newval.as_one())); +#endif +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { + rhs1 += rhs2; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + atomicAdd(&rhs1, rhs2); +#else +#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) + std::atomic_ref atomic(rhs1); + RHS oldval = atomic.load(); + RHS newval; + do { + newval = oldval + rhs2; + } while (!atomic.compare_exchange_weak(oldval, newval)); +#else + // No atomic floating point operations so use compare and swap + TypePunning::Alias oldval, newval; + TypePunning::Pointer pointer((void *)&rhs1); + do { + oldval.load(pointer); + newval = oldval.as_two() + rhs2; + } while (!__sync_bool_compare_and_swap( + (int32_t *)pointer, oldval.as_one(), newval.as_one())); +#endif +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { + lhs += rhs; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::apply(LHS &lhs, + RHS rhs) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) +#if (__CUDA_ARCH__ >= 600) || defined(__HIP_DEVICE_COMPILE__) + atomicAdd(&lhs, rhs); +#else + RHS newval = lhs, oldval; + // Type punning like this is illegal in C++ but the + // CUDA manual has an example just like it so fuck it + unsigned long long int *ptr = (unsigned long long int *)&lhs; + do { + oldval = newval; + newval += rhs; + newval = __ulonglong_as_double(atomicCAS( + ptr, __double_as_ulonglong(oldval), __double_as_ulonglong(newval))); + } while (oldval != newval); +#endif +#else +#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) + std::atomic_ref atomic(lhs); + RHS oldval = atomic.load(); + RHS newval; + do { + newval = oldval + rhs; + } while (!atomic.compare_exchange_weak(oldval, newval)); +#else + // No atomic floating point operations so use compare and swap + TypePunning::Alias oldval, newval; + TypePunning::Pointer pointer((void *)&lhs); + do { + oldval.load(pointer); + newval = oldval.as_two() + rhs; + } while (!__sync_bool_compare_and_swap( + (int64_t *)pointer, oldval.as_one(), newval.as_one())); +#endif +#endif +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { + rhs1 += rhs2; +} + +template <> +__LEGION_CUDA_HD__ inline void SumReduction::fold(RHS &rhs1, + RHS rhs2) { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) +#if (__CUDA_ARCH__ >= 600) || defined(__HIP_DEVICE_COMPILE__) + atomicAdd(&rhs1, rhs2); +#else + RHS newval = rhs1, oldval; + // Type punning like this is illegal in C++ but the + // CUDA manual has an example just like it so fuck it + unsigned long long int *ptr = (unsigned long long int *)&rhs1; + do { + oldval = newval; + newval += rhs2; + newval = __ulonglong_as_double(atomicCAS( + ptr, __double_as_ulonglong(oldval), __double_as_ulonglong(newval))); + } while (oldval != newval); +#endif +#else +#if defined(__cpp_lib_atomic_ref) && (__cpp_lib_atomic_ref >= 201806L) + std::atomic_ref atomic(rhs1); + RHS oldval = atomic.load(); + RHS newval; + do { + newval = oldval + rhs2; + } while (!atomic.compare_exchange_weak(oldval, newval)); +#else + // No atomic floating point operations so use compare and swap + TypePunning::Alias oldval, newval; + TypePunning::Pointer pointer((void *)&rhs1); + do { + oldval.load(pointer); + newval = oldval.as_two() + rhs2; + } while (!__sync_bool_compare_and_swap( + (int64_t *)pointer, oldval.as_one(), newval.as_one())); +#endif +#endif +} + +// Follow the Legion style of defining reduction operators first and then +// adapting the interface for CUDA +template +class AddCudaReductions : public T { +public: + static bool const has_cuda_reductions = true; + + template + static __device__ void apply_cuda(typename T::LHS &lhs, typename T::RHS rhs) { + T::template apply(lhs, rhs); + } + + template + static __device__ void fold_cuda(typename T::LHS &lhs, typename T::RHS rhs) { + T::template fold(lhs, rhs); + } +}; + +void register_all_redops(Realm::Runtime rt) { + // Registration is synchronous, so no need to capture events here + rt.register_reduction>>( + get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_BOOL_REDOP_ID)); + rt.register_reduction>>( + get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_INT32_REDOP_ID)); + rt.register_reduction>>( + get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_INT64_REDOP_ID)); + rt.register_reduction>>( + get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_FLOAT_REDOP_ID)); + rt.register_reduction>>( + get_realm_reduction_op_id_for_redop_id(redop_id_t::SUM_DOUBLE_REDOP_ID)); +} + +} // namespace FlexFlow diff --git a/lib/realm-execution/src/realm-execution/redops/redop_id_t.cc b/lib/realm-execution/src/realm-execution/redops/redop_id_t.cc index f31769419f..a4554c6f26 100644 --- a/lib/realm-execution/src/realm-execution/redops/redop_id_t.cc +++ b/lib/realm-execution/src/realm-execution/redops/redop_id_t.cc @@ -22,7 +22,7 @@ redop_id_t get_sum_redop_id_for_data_type(DataType dtype) { Realm::ReductionOpID get_realm_reduction_op_id_for_redop_id(redop_id_t redop_id) { - return static_cast(redop_id); + return static_cast(redop_id) + 1; } } // namespace FlexFlow