alpaka
Abstraction Library for Parallel Kernel Acceleration
Loading...
Searching...
No Matches
WorkDivGenericSycl.hpp
Go to the documentation of this file.
1/* Copyright 2023 Jan Stephan, Luca Ferragina, Andrea Bocci, Aurora Perego
2 * SPDX-License-Identifier: MPL-2.0
3 */
4
5#pragma once
6
8#include "alpaka/vec/Vec.hpp"
10
11#ifdef ALPAKA_ACC_SYCL_ENABLED
12
13# include <sycl/sycl.hpp>
14
15namespace alpaka
16{
17 //! The SYCL accelerator work division.
18 template<typename TDim, typename TIdx>
19 class WorkDivGenericSycl : public interface::Implements<ConceptWorkDiv, WorkDivGenericSycl<TDim, TIdx>>
20 {
21 static_assert(TDim::value > 0, "The SYCL work division must have a dimension greater than zero.");
22
23 public:
24 using WorkDivBase = WorkDivGenericSycl;
25
26 WorkDivGenericSycl(Vec<TDim, TIdx> const& threadElemExtent, sycl::nd_item<TDim::value> work_item)
27 : m_threadElemExtent{threadElemExtent}
28 , m_item_workdiv{work_item}
29 {
30 }
31
32 Vec<TDim, TIdx> const& m_threadElemExtent;
33 sycl::nd_item<TDim::value> m_item_workdiv;
34 };
35} // namespace alpaka
36
37namespace alpaka::trait
38{
39 //! The SYCL accelerator work division dimension get trait specialization.
40 template<typename TDim, typename TIdx>
41 struct DimType<WorkDivGenericSycl<TDim, TIdx>>
42 {
43 using type = TDim;
44 };
45
46 //! The SYCL accelerator work division idx type trait specialization.
47 template<typename TDim, typename TIdx>
48 struct IdxType<WorkDivGenericSycl<TDim, TIdx>>
49 {
50 using type = TIdx;
51 };
52
53 //! The SYCL accelerator work division grid block extent trait specialization.
54 template<typename TDim, typename TIdx>
55 struct GetWorkDiv<WorkDivGenericSycl<TDim, TIdx>, origin::Grid, unit::Blocks>
56 {
57 //! \return The number of blocks in each dimension of the grid.
58 static auto getWorkDiv(WorkDivGenericSycl<TDim, TIdx> const& workDiv) -> Vec<TDim, TIdx>
59 {
60 if constexpr(TDim::value == 0)
61 return Vec<TDim, TIdx>{};
62 else if constexpr(TDim::value == 1)
63 return Vec<TDim, TIdx>{static_cast<TIdx>(workDiv.m_item_workdiv.get_group_range(0))};
64 else if constexpr(TDim::value == 2)
65 {
66 return Vec<TDim, TIdx>{
67 static_cast<TIdx>(workDiv.m_item_workdiv.get_group_range(1)),
68 static_cast<TIdx>(workDiv.m_item_workdiv.get_group_range(0))};
69 }
70 else
71 {
72 return Vec<TDim, TIdx>{
73 static_cast<TIdx>(workDiv.m_item_workdiv.get_group_range(2)),
74 static_cast<TIdx>(workDiv.m_item_workdiv.get_group_range(1)),
75 static_cast<TIdx>(workDiv.m_item_workdiv.get_group_range(0))};
76 }
77 }
78 };
79
80 //! The SYCL accelerator work division block thread extent trait specialization.
81 template<typename TDim, typename TIdx>
82 struct GetWorkDiv<WorkDivGenericSycl<TDim, TIdx>, origin::Block, unit::Threads>
83 {
84 //! \return The number of threads in each dimension of a block.
85 static auto getWorkDiv(WorkDivGenericSycl<TDim, TIdx> const& workDiv) -> Vec<TDim, TIdx>
86 {
87 if constexpr(TDim::value == 0)
88 return Vec<TDim, TIdx>{};
89 else if constexpr(TDim::value == 1)
90 return Vec<TDim, TIdx>{static_cast<TIdx>(workDiv.m_item_workdiv.get_local_range(0))};
91 else if constexpr(TDim::value == 2)
92 {
93 return Vec<TDim, TIdx>{
94 static_cast<TIdx>(workDiv.m_item_workdiv.get_local_range(1)),
95 static_cast<TIdx>(workDiv.m_item_workdiv.get_local_range(0))};
96 }
97 else
98 {
99 return Vec<TDim, TIdx>{
100 static_cast<TIdx>(workDiv.m_item_workdiv.get_local_range(2)),
101 static_cast<TIdx>(workDiv.m_item_workdiv.get_local_range(1)),
102 static_cast<TIdx>(workDiv.m_item_workdiv.get_local_range(0))};
103 }
104 }
105 };
106
107 //! The SYCL accelerator work division thread element extent trait specialization.
108 template<typename TDim, typename TIdx>
109 struct GetWorkDiv<WorkDivGenericSycl<TDim, TIdx>, origin::Thread, unit::Elems>
110 {
111 //! \return The number of elements in each dimension of the thread.
112 static auto getWorkDiv(WorkDivGenericSycl<TDim, TIdx> const& workDiv) -> Vec<TDim, TIdx>
113 {
114 return workDiv.m_threadElemExtent;
115 }
116 };
117} // namespace alpaka::trait
118
119#endif
The accelerator traits.
The alpaka accelerator library.
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getWorkDiv(TWorkDiv const &workDiv) -> Vec< Dim< TWorkDiv >, Idx< TWorkDiv > >
Get the extent requested.
Definition Traits.hpp:33