alpaka
Abstraction Library for Parallel Kernel Acceleration
IdxBtUniformCudaHipBuiltIn.hpp
Go to the documentation of this file.
1 /* Copyright 2022 Axel Huebl, Benjamin Worpitz, Matthias Werner, RenĂ© Widera, Jan Stephan, Andrea Bocci, Bernhard
2  * Manfred Gruber
3  * SPDX-License-Identifier: MPL-2.0
4  */
5 
6 #pragma once
7 
10 #include "alpaka/core/Cuda.hpp"
11 #include "alpaka/core/Hip.hpp"
13 #include "alpaka/idx/Traits.hpp"
14 #include "alpaka/vec/Vec.hpp"
15 
16 #if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED)
17 
18 namespace alpaka
19 {
20  namespace bt
21  {
22  //! The CUDA/HIP accelerator ND index provider.
23  template<typename TDim, typename TIdx>
25  : public concepts::Implements<ConceptIdxBt, IdxBtUniformCudaHipBuiltIn<TDim, TIdx>>
26  {
27  };
28  } // namespace bt
29 
30 # if !defined(ALPAKA_HOST_ONLY)
31 
32 # if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) && !BOOST_LANG_CUDA
33 # error If ALPAKA_ACC_GPU_CUDA_ENABLED is set, the compiler has to support CUDA!
34 # endif
35 
36 # if defined(ALPAKA_ACC_GPU_HIP_ENABLED) && !BOOST_LANG_HIP
37 # error If ALPAKA_ACC_GPU_HIP_ENABLED is set, the compiler has to support HIP!
38 # endif
39 
40  namespace trait
41  {
42  //! The GPU CUDA/HIP accelerator index dimension get trait specialization.
43  template<typename TDim, typename TIdx>
44  struct DimType<bt::IdxBtUniformCudaHipBuiltIn<TDim, TIdx>>
45  {
46  using type = TDim;
47  };
48 
49  //! The GPU CUDA/HIP accelerator block thread index get trait specialization.
50  template<typename TDim, typename TIdx>
51  struct GetIdx<bt::IdxBtUniformCudaHipBuiltIn<TDim, TIdx>, origin::Block, unit::Threads>
52  {
53  //! \return The index of the current thread in the block.
54  template<typename TWorkDiv>
55  __device__ static auto getIdx(bt::IdxBtUniformCudaHipBuiltIn<TDim, TIdx> const& /* idx */, TWorkDiv const&)
57  {
58 # if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
59  return castVec<TIdx>(getOffsetVecEnd<TDim>(threadIdx));
60 # else
61  return getOffsetVecEnd<TDim>(Vec<std::integral_constant<typename TDim::value_type, 3>, TIdx>(
62  static_cast<TIdx>(hipThreadIdx_z),
63  static_cast<TIdx>(hipThreadIdx_y),
64  static_cast<TIdx>(hipThreadIdx_x)));
65 # endif
66  }
67  };
68 
69  //! The GPU CUDA/HIP accelerator block thread index idx type trait specialization.
70  template<typename TDim, typename TIdx>
71  struct IdxType<bt::IdxBtUniformCudaHipBuiltIn<TDim, TIdx>>
72  {
73  using type = TIdx;
74  };
75  } // namespace trait
76 
77 # endif
78 
79 } // namespace alpaka
80 
81 #endif
The CUDA/HIP accelerator ND index provider.
The alpaka accelerator library.
Tag used in class inheritance hierarchies that describes that a specific concept (TConcept) is implem...
Definition: Concepts.hpp:15
The dimension getter type trait.
Definition: Traits.hpp:14
static __device__ auto getIdx(bt::IdxBtUniformCudaHipBuiltIn< TDim, TIdx > const &, TWorkDiv const &) -> Vec< TDim, TIdx >
The index get trait.
Definition: Traits.hpp:42
The idx type trait.
Definition: Traits.hpp:25