AccGenericSycl.hpp

Go to the documentation of this file.

/* Copyright 2026 Jan Stephan, Antonio Di Pilato, Andrea Bocci, Luca Ferragina, Aurora Perego
 * SPDX-License-Identifier: MPL-2.0
 */
 
#pragma once
 
// Base classes.
#include "alpaka/atomic/AtomicGenericSycl.hpp"
#include "alpaka/atomic/AtomicHierarchy.hpp"
#include "alpaka/block/shared/dyn/BlockSharedMemDynGenericSycl.hpp"
#include "alpaka/block/shared/st/BlockSharedMemStGenericSycl.hpp"
#include "alpaka/block/sync/BlockSyncGenericSycl.hpp"
#include "alpaka/dev/DevGenericSycl.hpp"
#include "alpaka/grid/GridSyncGenericSycl.hpp"
#include "alpaka/idx/bt/IdxBtGenericSycl.hpp"
#include "alpaka/idx/gb/IdxGbGenericSycl.hpp"
#include "alpaka/intrinsic/IntrinsicGenericSycl.hpp"
#include "alpaka/math/MathGenericSycl.hpp"
#include "alpaka/mem/fence/MemFenceGenericSycl.hpp"
#include "alpaka/platform/PlatformGenericSycl.hpp"
#include "alpaka/rand/RandDefault.hpp"
#include "alpaka/rand/RandGenericSycl.hpp"
#include "alpaka/warp/WarpGenericSycl.hpp"
#include "alpaka/workdiv/WorkDivGenericSycl.hpp"
 
// Specialized traits.
#include "alpaka/acc/Traits.hpp"
#include "alpaka/dev/Traits.hpp"
#include "alpaka/idx/Traits.hpp"
#include "alpaka/kernel/Traits.hpp"
#include "alpaka/platform/Traits.hpp"
#include "alpaka/vec/Vec.hpp"
 
// Implementation details.
#include "alpaka/core/ClipCast.hpp"
#include "alpaka/core/Config.hpp"
#include "alpaka/core/Interface.hpp"
#include "alpaka/core/Sycl.hpp"
 
#include <algorithm>
#include <cstddef>
#ifdef __cpp_lib_format
#    include <format>
#endif
#include <string>
#include <type_traits>
 
#ifdef ALPAKA_ACC_SYCL_ENABLED
 
#    include <sycl/sycl.hpp>
 
namespace alpaka
{
    template<
        concepts::Tag TTag,
        typename TAcc,
        typename TDim,
        typename TIdx,
        typename TKernelFnObj,
        bool TCooperative,
        typename... TArgs>
    class TaskKernelGenericSycl;
 
    //! The SYCL accelerator.
    //!
    //! This accelerator allows parallel kernel execution on SYCL devices.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    class AccGenericSycl
        : public WorkDivGenericSycl<TDim, TIdx>
        , public gb::IdxGbGenericSycl<TDim, TIdx>
        , public bt::IdxBtGenericSycl<TDim, TIdx>
        , public AtomicHierarchy<AtomicGenericSycl, AtomicGenericSycl, AtomicGenericSycl>
        , public math::MathGenericSycl
        , public BlockSharedMemDynGenericSycl
        , public BlockSharedMemStGenericSycl
        , public BlockSyncGenericSycl<TDim>
        , public GridSyncGenericSycl<TDim>
        , public IntrinsicGenericSycl
        , public MemFenceGenericSycl
#    ifdef ALPAKA_DISABLE_VENDOR_RNG
        , public rand::RandDefault
#    else
        , public rand::RandGenericSycl<TDim>
#    endif
        , public warp::WarpGenericSycl<TDim>
        , public interface::Implements<InterfaceAcc, AccGenericSycl<TTag, TDim, TIdx>>
    {
        static_assert(TDim::value > 0, "The SYCL accelerator must have a dimension greater than zero.");
 
    public:
        AccGenericSycl(AccGenericSycl const&) = delete;
        AccGenericSycl(AccGenericSycl&&) = delete;
        auto operator=(AccGenericSycl const&) -> AccGenericSycl& = delete;
        auto operator=(AccGenericSycl&&) -> AccGenericSycl& = delete;
 
        AccGenericSycl(
            Vec<TDim, TIdx> const& threadElemExtent,
            sycl::nd_item<TDim::value> work_item,
            sycl::local_accessor<std::byte> dyn_shared_acc,
            sycl::local_accessor<std::byte> st_shared_acc)
            : WorkDivGenericSycl<TDim, TIdx>{threadElemExtent, work_item}
            , gb::IdxGbGenericSycl<TDim, TIdx>{work_item}
            , bt::IdxBtGenericSycl<TDim, TIdx>{work_item}
            , BlockSharedMemDynGenericSycl{dyn_shared_acc}
            , BlockSharedMemStGenericSycl{st_shared_acc}
            , BlockSyncGenericSycl<TDim>{work_item}
            , GridSyncGenericSycl<TDim>{work_item}
#    ifndef ALPAKA_DISABLE_VENDOR_RNG
            , rand::RandGenericSycl<TDim>{work_item}
#    endif
            , warp::WarpGenericSycl<TDim>{work_item}
        {
        }
    };
} // namespace alpaka
 
namespace alpaka::detail
{
    //! Whether the SYCL device can run cooperative kernels, i.e. kernels that synchronise all the work-items in the
    //! grid: this requires the work-groups in the root group to make concurrent forward progress.
    template<concepts::Tag TTag>
    auto syclCooperativeLaunch([[maybe_unused]] sycl::device const& device) -> bool
    {
        if constexpr(std::is_same_v<TTag, TagFpgaSyclIntel>)
        {
            // cooperative kernels are not supported by the FPGA back-end
            return false;
        }
        else
        {
#    ifdef SYCL_EXT_ONEAPI_FORWARD_PROGRESS
            namespace syclex = sycl::ext::oneapi::experimental;
            auto const caps = device.template get_info<
                syclex::info::device::work_group_progress_capabilities<syclex::execution_scope::root_group>>();
            return std::find(caps.begin(), caps.end(), syclex::forward_progress_guarantee::concurrent) != caps.end();
#    else
            // without the forward progress extension there is no way to check if the device supports it
            return false;
#    endif
        }
    }
} // namespace alpaka::detail
 
namespace alpaka::trait
{
    //! The SYCL accelerator type trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct AccType<AccGenericSycl<TTag, TDim, TIdx>>
    {
        using type = AccGenericSycl<TTag, TDim, TIdx>;
    };
 
    //! The SYCL single thread accelerator type trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct IsSingleThreadAcc<AccGenericSycl<TTag, TDim, TIdx>> : std::false_type
    {
    };
 
    //! The SYCL multi thread accelerator type trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct IsMultiThreadAcc<AccGenericSycl<TTag, TDim, TIdx>> : std::true_type
    {
    };
 
    //! The SYCL accelerator device properties get trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct GetAccDevProps<AccGenericSycl<TTag, TDim, TIdx>>
    {
        static auto getAccDevProps(DevGenericSycl<TTag> const& dev) -> AccDevProps<TDim, TIdx>
        {
            auto const device = dev.getNativeHandle().first;
            auto const max_threads_dim
                = device.template get_info<sycl::info::device::max_work_item_sizes<TDim::value>>();
            Vec<TDim, TIdx> max_threads_dim_vec{};
            for(int i = 0; i < static_cast<int>(TDim::value); i++)
                max_threads_dim_vec[i] = alpaka::core::clipCast<TIdx>(max_threads_dim[i]);
            return {// m_multiProcessorCount
                    alpaka::core::clipCast<TIdx>(device.template get_info<sycl::info::device::max_compute_units>()),
                    // m_gridBlockExtentMax
                    getExtentVecEnd<TDim>(Vec<DimInt<3u>, TIdx>(
                        // WARNING: There is no SYCL way to determine these values
                        std::numeric_limits<TIdx>::max(),
                        std::numeric_limits<TIdx>::max(),
                        std::numeric_limits<TIdx>::max())),
                    // m_gridBlockCountMax
                    std::numeric_limits<TIdx>::max(),
                    // m_blockThreadExtentMax
                    max_threads_dim_vec,
                    // m_blockThreadCountMax
                    alpaka::core::clipCast<TIdx>(device.template get_info<sycl::info::device::max_work_group_size>()),
                    // m_threadElemExtentMax
                    Vec<TDim, TIdx>::all(std::numeric_limits<TIdx>::max()),
                    // m_threadElemCountMax
                    std::numeric_limits<TIdx>::max(),
                    // m_sharedMemSizeBytes
                    device.template get_info<sycl::info::device::local_mem_size>(),
                    // m_globalMemSizeBytes
                    getMemBytes(dev),
                    // m_cooperativeLaunch
                    alpaka::detail::syclCooperativeLaunch<TTag>(device)};
        }
    };
 
    //! The SYCL accelerator name trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct GetAccName<AccGenericSycl<TTag, TDim, TIdx>>
    {
        static auto getAccName() -> std::string
        {
#    if ALPAKA_COMP_CLANG
#        pragma clang diagnostic push
#        pragma clang diagnostic ignored "-Wexit-time-destructors"
#    endif
            using namespace std::literals;
            static std::string const accName =
#    ifdef __cpp_lib_format
                std::format(
                    "Acc{}<{},{}>",
                    core::demangled<TTag>.substr(std::string_view("alpaka::Tag").size()),
                    TDim::value,
                    core::demangled<TIdx>);
#    else
                "Acc"s + std::string(core::demangled<TTag>.substr(std::string_view("alpaka::Tag").size())) + "<"s
                + std::to_string(TDim::value) + ","s + std::string(core::demangled<TIdx>) + ">"s;
#    endif
            return accName;
#    if ALPAKA_COMP_CLANG
#        pragma clang diagnostic pop
#    endif
        }
    };
 
    //! The SYCL accelerator device type trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct DevType<AccGenericSycl<TTag, TDim, TIdx>>
    {
        using type = DevGenericSycl<TTag>;
    };
 
    //! The SYCL accelerator dimension getter trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct DimType<AccGenericSycl<TTag, TDim, TIdx>>
    {
        using type = TDim;
    };
 
    //! The SYCL accelerator execution task type trait specialization.
    template<
        concepts::Tag TTag,
        typename TDim,
        typename TIdx,
        typename TWorkDiv,
        typename TKernelFnObj,
        typename... TArgs>
    struct CreateTaskKernel<AccGenericSycl<TTag, TDim, TIdx>, TWorkDiv, TKernelFnObj, TArgs...>
    {
        static auto createTaskKernel(TWorkDiv const& workDiv, TKernelFnObj const& kernelFnObj, TArgs&&... args)
        {
            return TaskKernelGenericSycl<
                TTag,
                AccGenericSycl<TTag, TDim, TIdx>,
                TDim,
                TIdx,
                TKernelFnObj,
                false,
                TArgs...>{workDiv, kernelFnObj, std::forward<TArgs>(args)...};
        }
    };
 
    //! The SYCL accelerator execution task type trait specialization.
    template<typename TTag, typename TDim, typename TIdx, typename TWorkDiv, typename TKernelFnObj, typename... TArgs>
    struct CreateTaskCooperativeKernel<AccGenericSycl<TTag, TDim, TIdx>, TWorkDiv, TKernelFnObj, TArgs...>
    {
        static auto createTaskCooperativeKernel(
            TWorkDiv const& workDiv,
            TKernelFnObj const& kernelFnObj,
            TArgs&&... args)
        {
            return TaskKernelGenericSycl<
                TTag,
                AccGenericSycl<TTag, TDim, TIdx>,
                TDim,
                TIdx,
                TKernelFnObj,
                true,
                TArgs...>{workDiv, kernelFnObj, std::forward<TArgs>(args)...};
        }
    };
 
    //! The SYCL execution task platform type trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct PlatformType<AccGenericSycl<TTag, TDim, TIdx>>
    {
        using type = PlatformGenericSycl<TTag>;
    };
 
    //! The SYCL accelerator idx type trait specialization.
    template<concepts::Tag TTag, typename TDim, typename TIdx>
    struct IdxType<AccGenericSycl<TTag, TDim, TIdx>>
    {
        using type = TIdx;
    };
} // namespace alpaka::trait
 
#endif