group.hpp

Go to the documentation of this file.

#ifndef TDLS_CORE_GROUP_HPP
#define TDLS_CORE_GROUP_HPP
 
 
 
 
 
 
#include <atomic>
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <thread>
#include <type_traits>
#include <utility>
 
#include <tdls/core/macros.hpp>
#include <tdls/core/target.hpp>
 
#if defined(TDLS_DETAIL_COMPILER_NVCXX)
#include <nv/target>
#endif
 
// The device printf of HIP comes with its runtime header, which a HIP
// translation unit may include after this one.
#if defined(TDLS_DETAIL_COMPILER_HIP_CLANG)
#include <hip/hip_runtime.h>
#endif
 
 
 
/* =========================================================================
   Side of the compilation pass (TDLS_DETAIL_PASS_*, core/target.hpp): the
   instructions it offers to run the deduced barrier and to stop a kernel
   or the program. The NVIDIA builtins serve the device passes outside the
   CUDA language only, and the AMD ones the device passes of HIP and of
   the other models. A pass without a side has no deduced barrier, and
   TDLS_DETAIL_GROUP_NONE gives the reason.
   ========================================================================= */
 
#if defined(TDLS_DETAIL_PASS_NVCXX_STDPAR_OPENACC)
#define TDLS_DETAIL_GROUP_NONE "The OpenACC backend of nvc++ for the parallel algorithms has none."
#elif defined(TDLS_DETAIL_PASS_NVCXX_CUDA)
#define TDLS_DETAIL_GROUP_NVCXX
#elif defined(TDLS_DETAIL_PASS_NVCXX_OFFLOAD)
#define TDLS_DETAIL_GROUP_NONE "nvc++ offers it to OpenMP and OpenACC offloading with -cuda."
#elif defined(TDLS_DETAIL_PASS_AMDGCN)
#define TDLS_DETAIL_GROUP_AMDGCN
#elif defined(TDLS_DETAIL_PASS_CUDA)
#define TDLS_DETAIL_GROUP_CUDA
#elif defined(TDLS_DETAIL_PASS_NVPTX)
#define TDLS_DETAIL_GROUP_NVPTX
#elif defined(TDLS_DETAIL_PASS_GCC_OFFLOAD)
#define TDLS_DETAIL_GROUP_NONE "GCC offloading has none."
#elif defined(TDLS_DETAIL_PASS_SPIRV)
#define TDLS_DETAIL_GROUP_NONE                                                                     \
    "On SPIR-V devices, SYCL and OpenMP offloading cannot stop a kernel whose group is "           \
    "incomplete; pass a sub-group, work-group or team barrier."
#elif defined(TDLS_DETAIL_PASS_ACPP_GENERIC)
#define TDLS_DETAIL_GROUP_NONE                                                                     \
    "The generic mode of AdaptiveCpp has none; its cuda and hip modes do."
#else
#define TDLS_DETAIL_GROUP_HOST
#endif
 
// Two sides lack the deduced barrier, yet still stop a kernel or the
// program: NVIDIA before Volta has no match instruction, and a standard
// library without std::atomic_ref no CPU barrier on the workspace.
#if (defined(TDLS_DETAIL_GROUP_CUDA) || defined(TDLS_DETAIL_GROUP_NVPTX)) &&                       \
    ((defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700) ||                                            \
     (defined(__SYCL_CUDA_ARCH__) && __SYCL_CUDA_ARCH__ < 700))
#define TDLS_DETAIL_GROUP_NONE "NVIDIA GPUs before Volta (sm_70) have none."
#elif defined(TDLS_DETAIL_GROUP_HOST) && !defined(__cpp_lib_atomic_ref)
#define TDLS_DETAIL_GROUP_NONE                                                                     \
    "On CPU, it needs std::atomic_ref, missing from this standard library (libc++ before 19)."
#endif
 
#if !defined(TDLS_DETAIL_GROUP_NONE)
#define TDLS_DETAIL_GROUP_BARRIER
#endif
 
#define TDLS_DETAIL_GROUP_EXPAND(...) __VA_ARGS__
 
// The arguments of the selection macros are parenthesized statement
// lists, as for NV_IF_TARGET. TDLS_DETAIL_GROUP_LANES qualifies the lane
// functions: device functions of the CUDA and HIP languages, host-device
// functions of nvc++, which select their body with NV_IF_TARGET, and
// plain functions of the device passes of the other models.
#if defined(TDLS_DETAIL_GROUP_NVCXX)
#define TDLS_DETAIL_GROUP_ON_DEVICE(code) NV_IF_TARGET(NV_IS_DEVICE, code)
#define TDLS_DETAIL_GROUP_SIDE(device, host) NV_IF_TARGET(NV_IS_DEVICE, device, host)
#define TDLS_DETAIL_GROUP_LANES TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU
#elif defined(TDLS_DETAIL_GROUP_CUDA)
#define TDLS_DETAIL_GROUP_ON_DEVICE(code)    TDLS_DETAIL_GROUP_EXPAND code
#define TDLS_DETAIL_GROUP_SIDE(device, host) TDLS_DETAIL_GROUP_EXPAND device
#define TDLS_DETAIL_GROUP_LANES              __device__ TDLS_FORCE_INLINE_GPU
#elif defined(TDLS_DETAIL_GROUP_AMDGCN) && defined(TDLS_DETAIL_COMPILER_HIP_CLANG)
#define TDLS_DETAIL_GROUP_SIDE(device, host) TDLS_DETAIL_GROUP_EXPAND device
#define TDLS_DETAIL_GROUP_LANES              __device__ TDLS_FORCE_INLINE_GPU
#elif defined(TDLS_DETAIL_GROUP_NVPTX) || defined(TDLS_DETAIL_GROUP_AMDGCN)
#define TDLS_DETAIL_GROUP_SIDE(device, host) TDLS_DETAIL_GROUP_EXPAND device
#define TDLS_DETAIL_GROUP_LANES              inline
#elif defined(TDLS_DETAIL_GROUP_HOST)
#define TDLS_DETAIL_GROUP_SIDE(device, host) TDLS_DETAIL_GROUP_EXPAND host
#else
#define TDLS_DETAIL_GROUP_SIDE(device, host)
#endif
 
#if defined(TDLS_DETAIL_GROUP_BARRIER)
#define TDLS_DETAIL_GROUP_DISPATCH(device, host) TDLS_DETAIL_GROUP_SIDE(device, host)
#else
#define TDLS_DETAIL_GROUP_DISPATCH(device, host)
#endif
 
 
 
namespace tdls {
 
 
 
struct AutoSync {};
 
struct NoSync {
    TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void operator()() const noexcept {
    }
};
 
namespace detail {
 
template<typename Sync>
inline constexpr bool is_auto_sync = std::is_same_v<std::remove_cvref_t<Sync>, AutoSync>;
 
template<typename Sync>
inline constexpr bool is_no_sync = std::is_same_v<std::remove_cvref_t<Sync>, NoSync>;
 
template<typename Sync, bool = is_auto_sync<Sync>>
inline constexpr bool nothrow_sync = noexcept(std::declval<Sync&>()());
 
template<typename Sync>
inline constexpr bool nothrow_sync<Sync, true> = true;
 
#if defined(TDLS_DETAIL_GROUP_BARRIER)
template<typename T>
inline constexpr bool group_barrier_available = true;
#define TDLS_DETAIL_GROUP_NONE ""
#else
template<typename T>
inline constexpr bool group_barrier_available = false;
#endif
 
inline constexpr int barrier_elements = 2;
 
#if defined(TDLS_DETAIL_GROUP_BARRIER) &&                                                          \
    (defined(TDLS_DETAIL_GROUP_CUDA) || defined(TDLS_DETAIL_GROUP_NVPTX) ||                        \
     defined(TDLS_DETAIL_GROUP_NVCXX))
inline constexpr int group_lanes = 32;
#elif defined(TDLS_DETAIL_GROUP_BARRIER) && defined(TDLS_DETAIL_GROUP_AMDGCN)
inline constexpr int group_lanes = 64;
#else
inline constexpr int group_lanes = 0;
#endif
 
enum class LanesFailure {
    incomplete, 
    no_sync     
};
 
// The messages, printed by the first lane of the group. They take no
// argument: a formatted print would reserve a stack frame in local
// memory for every kernel, even when it never runs. A workspace in the
// private memory of each thread fails as an incomplete group: its address
// is the same in every lane, so the lanes would match without sharing
// anything.
#define TDLS_DETAIL_GROUP_INCOMPLETE                                                               \
    "TDLS: a group of threads is incomplete in its warp, or its workspace is private to each "     \
    "thread. The threads of a group must call the solver together, as lanes of one warp or "       \
    "wavefront, with a workspace in shared or global memory.\n"
#define TDLS_DETAIL_GROUP_NO_SYNC                                                                  \
    "TDLS: the runtime solver runs with tdls::NoSync, valid for one thread per system, on a "      \
    "dimension above rows_per_thread. Pass a barrier, or a rows_per_thread reaching the "          \
    "dimension.\n"
 
/* =========================================================================
   NVIDIA, CUDA API: nvcc, clang CUDA, nvc++ on the device side.
   ========================================================================= */
 
#if defined(TDLS_DETAIL_GROUP_CUDA) || defined(TDLS_DETAIL_GROUP_NVCXX)
 
TDLS_DETAIL_GROUP_LANES void lanes_fail([[maybe_unused]] const unsigned long long mask,
                                        [[maybe_unused]] const LanesFailure failure) noexcept {
    // clang-format off: a statement list as the argument of a macro
    TDLS_DETAIL_GROUP_ON_DEVICE((
        unsigned lane;
        asm volatile("mov.u32 %0, %%laneid;" : "=r"(lane));
        if (lane == static_cast<unsigned>(__ffsll(static_cast<long long>(mask)) - 1)) {
            if (failure == LanesFailure::incomplete)
                printf(TDLS_DETAIL_GROUP_INCOMPLETE);
            else
                printf(TDLS_DETAIL_GROUP_NO_SYNC);
        }
        __trap();))
    // clang-format on
}
 
TDLS_DETAIL_GROUP_LANES unsigned long long lanes_active() noexcept {
    unsigned long long lanes = 0;
    TDLS_DETAIL_GROUP_ON_DEVICE((lanes = __activemask();))
    return lanes;
}
 
#if defined(TDLS_DETAIL_GROUP_BARRIER)
 
TDLS_DETAIL_GROUP_LANES unsigned long long lanes_join(const void* key, const int threads) noexcept {
    unsigned long long lanes = 0;
    // clang-format off: a statement list as the argument of a macro
    TDLS_DETAIL_GROUP_ON_DEVICE((
        const unsigned mask =
            __match_any_sync(__activemask(), reinterpret_cast<unsigned long long>(key));
        if (__popc(mask) != threads || __isLocal(key))
            lanes_fail(mask, LanesFailure::incomplete);
        lanes = mask;))
    // clang-format on
    return lanes;
}
 
TDLS_DETAIL_GROUP_LANES void lanes_sync(const unsigned long long lanes) noexcept {
    TDLS_DETAIL_GROUP_ON_DEVICE((__syncwarp(static_cast<unsigned>(lanes));))
}
 
#endif
 
/* =========================================================================
   NVIDIA, clang builtins: OpenMP offloading and SYCL device passes. The
   active mask is read with the PTX instruction itself, as clang's CUDA
   headers do before __nvvm_activemask (LLVM 19).
   ========================================================================= */
 
#elif defined(TDLS_DETAIL_GROUP_NVPTX)
 
TDLS_DETAIL_GROUP_LANES void lanes_fail([[maybe_unused]] const unsigned long long mask,
                                        [[maybe_unused]] const LanesFailure failure) noexcept {
#if !defined(TDLS_DETAIL_LANG_SYCL_DEVICE)
    if (__nvvm_read_ptx_sreg_laneid() == __builtin_ctzll(mask)) {
        if (failure == LanesFailure::incomplete)
            printf(TDLS_DETAIL_GROUP_INCOMPLETE);
        else
            printf(TDLS_DETAIL_GROUP_NO_SYNC);
    }
#endif
    __builtin_trap();
}
 
TDLS_DETAIL_GROUP_LANES unsigned long long lanes_active() noexcept {
    unsigned active;
    asm volatile("activemask.b32 %0;" : "=r"(active));
    return active;
}
 
#if defined(TDLS_DETAIL_GROUP_BARRIER)
 
TDLS_DETAIL_GROUP_LANES unsigned long long lanes_join(const void* key, const int threads) noexcept {
    const unsigned mask = __nvvm_match_any_sync_i64(
        static_cast<unsigned>(lanes_active()),
        static_cast<long long>(reinterpret_cast<unsigned long long>(key)));
    if (__builtin_popcount(mask) != threads || __nvvm_isspacep_local(key))
        lanes_fail(mask, LanesFailure::incomplete);
    return mask;
}
 
TDLS_DETAIL_GROUP_LANES void lanes_sync(const unsigned long long lanes) noexcept {
    __nvvm_bar_warp_sync(static_cast<unsigned>(lanes));
}
 
#endif
 
/* =========================================================================
   AMD, clang builtins: HIP, OpenMP offloading and SYCL device passes.
   The match is the loop of the __match_any of HIP: the first lane still
   searching broadcasts its key, and the lanes holding it leave with
   their mask. The lanes of a wavefront run in lockstep: the barrier only
   orders the memory operations, as the __syncwarp of HIP. These are the
   builtins HIP itself uses, available from ROCm 5.3 on, so that the
   barrier does not depend on the HIP release. The ballot is the builtin
   of recent clang, or the one HIP used before it.
   ========================================================================= */
 
#elif defined(TDLS_DETAIL_GROUP_AMDGCN)
 
#if __clang_major__ >= 19
#define TDLS_DETAIL_GROUP_BALLOT(predicate) __builtin_amdgcn_ballot_w64(predicate)
#else
#define TDLS_DETAIL_GROUP_BALLOT(predicate)                                                        \
    __builtin_amdgcn_uicmp(static_cast<unsigned>(predicate), 0u, 33 /* ICMP_NE */)
#endif
 
TDLS_DETAIL_GROUP_LANES void lanes_fail([[maybe_unused]] const unsigned long long mask,
                                        [[maybe_unused]] const LanesFailure failure) noexcept {
#if !defined(TDLS_DETAIL_LANG_SYCL_DEVICE)
    const unsigned lane = __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
    if (lane == static_cast<unsigned>(__builtin_ctzll(mask))) {
        if (failure == LanesFailure::incomplete)
            printf(TDLS_DETAIL_GROUP_INCOMPLETE);
        else
            printf(TDLS_DETAIL_GROUP_NO_SYNC);
    }
#endif
    __builtin_trap();
}
 
TDLS_DETAIL_GROUP_LANES unsigned long long lanes_active() noexcept {
    return __builtin_amdgcn_read_exec();
}
 
TDLS_DETAIL_GROUP_LANES unsigned long long lanes_join(const void* key, const int threads) noexcept {
    const unsigned long long own = reinterpret_cast<unsigned long long>(key);
    unsigned long long mask      = 0;
    bool done                    = false;
    while (TDLS_DETAIL_GROUP_BALLOT(!done) != 0) {
        if (!done) {
            const unsigned low  = __builtin_amdgcn_readfirstlane(static_cast<unsigned>(own));
            const unsigned high = __builtin_amdgcn_readfirstlane(static_cast<unsigned>(own >> 32));
            if (((static_cast<unsigned long long>(high) << 32) | low) == own) {
                mask = TDLS_DETAIL_GROUP_BALLOT(true);
                done = true;
            }
        }
    }
    // The builtin takes a generic pointer, which OpenMP offloading does not
    // make the default. Only a C cast changes the address space, as in the
    // OpenMP device runtime.
    const bool isprivate =
        __builtin_amdgcn_is_private((const __attribute__((address_space(0))) void*)key);
    if (__builtin_popcountll(mask) != threads || isprivate)
        lanes_fail(mask, LanesFailure::incomplete);
    return mask;
}
 
TDLS_DETAIL_GROUP_LANES void lanes_sync(unsigned long long) noexcept {
    __builtin_amdgcn_fence(__ATOMIC_RELEASE, "wavefront");
    __builtin_amdgcn_wave_barrier();
    __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "wavefront");
}
 
#endif
 
/* =========================================================================
   Host: a counter and a flag in the first two elements of the workspace.
   The flag alternates between 0 and 1. A thread reads it before
   arriving, so it cannot miss the flip of its own barrier; the last
   thread to arrive resets the counter, then flips the flag.
   ========================================================================= */
 
#if defined(TDLS_DETAIL_GROUP_HOST) || defined(TDLS_DETAIL_GROUP_NVCXX)
 
inline void host_refuse_no_sync() noexcept {
    std::fputs(TDLS_DETAIL_GROUP_NO_SYNC, stderr);
    std::abort();
}
 
#if defined(TDLS_DETAIL_GROUP_BARRIER)
 
template<typename T>
void host_join(T* slots, const int threads) noexcept {
    static_assert(std::atomic_ref<T>::is_always_lock_free,
                  "TDLS: the deduced CPU barrier needs lock-free atomics on the scalar type; "
                  "pass a barrier (sync argument)");
    const T count = std::atomic_ref<T>(slots[0]).load(std::memory_order_relaxed);
    const T flag  = std::atomic_ref<T>(slots[1]).load(std::memory_order_relaxed);
    if (!(count >= T(0) && count < T(threads) && count == T(static_cast<int>(count)) &&
          (flag == T(0) || flag == T(1)))) {
        std::fprintf(stderr,
                     "TDLS: the workspace of a group of %d CPU threads was not zero before "
                     "its first use.\n",
                     threads);
        std::abort();
    }
}
 
template<typename T>
void host_sync(T* slots, const int threads) noexcept {
    std::atomic_ref<T> count(slots[0]);
    std::atomic_ref<T> flag(slots[1]);
    const T seen = flag.load(std::memory_order_acquire);
    if (count.fetch_add(T(1), std::memory_order_acq_rel) == T(threads - 1)) {
        count.store(T(0), std::memory_order_relaxed);
        flag.store(seen == T(0) ? T(1) : T(0), std::memory_order_release);
        return;
    }
    // Spin first, then yield. The clock is read only after the spins, once
    // every 1024 iterations: the barriers of a running group never reach it.
    static std::atomic<bool> warned{false};
    std::chrono::steady_clock::time_point start{};
    for (unsigned spins = 0; flag.load(std::memory_order_acquire) == seen; ++spins) {
        if (spins < 1024) continue;
        std::this_thread::yield();
        if (spins % 1024 != 0) continue;
        const auto now = std::chrono::steady_clock::now();
        if (spins == 1024)
            start = now;
        else if (now - start > std::chrono::seconds(1) && !warned.exchange(true))
            std::fprintf(stderr,
                         "TDLS: a group of %d CPU threads has waited for 1 s at a barrier. "
                         "The threads of a group must share its workspace and run "
                         "concurrently: not in a worksharing loop, nor in parallel STL or a "
                         "SYCL kernel on the CPU.\n",
                         threads);
    }
}
 
#endif
 
#endif
 
#undef TDLS_DETAIL_GROUP_INCOMPLETE
#undef TDLS_DETAIL_GROUP_NO_SYNC
#undef TDLS_DETAIL_GROUP_BALLOT
 
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU void refuse_no_sync() noexcept {
    TDLS_DETAIL_GROUP_SIDE((lanes_fail(lanes_active(), LanesFailure::no_sync);),
                           (host_refuse_no_sync();))
}
 
template<typename T>
class DeducedSync {
  public:
    TDLS_EXEC_CHECK_DISABLE
    TDLS_HOST_DEVICE
    TDLS_FORCE_INLINE_GPU constexpr DeducedSync(T* slots, const int threads) noexcept
        : slots(slots), threads(threads) {
        if (threads > 1) {
            TDLS_DETAIL_GROUP_DISPATCH((lanes = lanes_join(slots, threads);),
                                       (host_join(slots, threads);))
        }
    }
 
    TDLS_EXEC_CHECK_DISABLE
    TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr void operator()() const noexcept {
        if (threads > 1) {
            TDLS_DETAIL_GROUP_DISPATCH((lanes_sync(lanes);), (host_sync(slots, threads);))
        }
    }
 
  private:
    T* slots;
    int threads;
    unsigned long long lanes = 0;
};
 
template<int static_threads, typename Sync, typename T>
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr decltype(auto)
resolve_sync(Sync& sync, [[maybe_unused]] T* slots, [[maybe_unused]] const int threads) noexcept {
    if constexpr (is_no_sync<Sync>) {
        static_assert(static_threads <= 1, "TDLS: a group of several threads cannot run with "
                                           "tdls::NoSync (sync argument)");
        if constexpr (static_threads == 0) {
            if (threads > 1) refuse_no_sync();
        }
        return (sync);
    } else if constexpr (!is_auto_sync<Sync>) {
        return (sync);
    } else if constexpr (static_threads == 1) {
        return NoSync{};
    } else {
        static_assert(group_barrier_available<T>,
                      "TDLS: this compilation target offers no deduced barrier for a group of "
                      "several threads; pass a barrier (sync argument), or tdls::NoSync with one "
                      "thread per system. " TDLS_DETAIL_GROUP_NONE);
        static_assert(group_lanes == 0 || static_threads <= group_lanes,
                      "TDLS: a group larger than a warp needs a barrier (sync argument)");
        return DeducedSync<T>(slots, threads);
    }
}
 
} // namespace detail
 
 
 
} // namespace tdls
 
 
 
#undef TDLS_DETAIL_GROUP_DISPATCH
#undef TDLS_DETAIL_GROUP_SIDE
#undef TDLS_DETAIL_GROUP_LANES
#undef TDLS_DETAIL_GROUP_ON_DEVICE
#undef TDLS_DETAIL_GROUP_EXPAND
#undef TDLS_DETAIL_GROUP_BARRIER
#undef TDLS_DETAIL_GROUP_CUDA
#undef TDLS_DETAIL_GROUP_NVPTX
#undef TDLS_DETAIL_GROUP_AMDGCN
#undef TDLS_DETAIL_GROUP_NVCXX
#undef TDLS_DETAIL_GROUP_HOST
#undef TDLS_DETAIL_GROUP_NONE
 
 
 
#endif // TDLS_CORE_GROUP_HPP