alpaka
Abstraction Library for Parallel Kernel Acceleration
BlockSharedMemStUniformCudaHipBuiltIn.hpp
Go to the documentation of this file.
1 /* Copyright 2022 Benjamin Worpitz, Erik Zenker, RenĂ© Widera, Matthias Werner, Andrea Bocci, Bernhard Manfred Gruber
2  * SPDX-License-Identifier: MPL-2.0
3  */
4 
5 #pragma once
6 
10 
11 #include <cstdint>
12 #include <type_traits>
13 
14 #if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED)
15 
16 namespace alpaka
17 {
18  //! The GPU CUDA/HIP block shared memory allocator.
20  : public concepts::Implements<ConceptBlockSharedSt, BlockSharedMemStUniformCudaHipBuiltIn>
21  {
22  };
23 
24 # if !defined(ALPAKA_HOST_ONLY)
25 
26 # if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) && !BOOST_LANG_CUDA
27 # error If ALPAKA_ACC_GPU_CUDA_ENABLED is set, the compiler has to support CUDA!
28 # endif
29 
30 # if defined(ALPAKA_ACC_GPU_HIP_ENABLED) && !BOOST_LANG_HIP
31 # error If ALPAKA_ACC_GPU_HIP_ENABLED is set, the compiler has to support HIP!
32 # endif
33 
34  namespace trait
35  {
36  template<typename T, std::size_t TuniqueId>
38  {
39  __device__ static auto declareVar(BlockSharedMemStUniformCudaHipBuiltIn const&) -> T&
40  {
41  __shared__ uint8_t shMem alignas(alignof(T))[sizeof(T)];
42  return *(reinterpret_cast<T*>(shMem));
43  }
44  };
45 
46  template<>
48  {
49  __device__ static auto freeVars(BlockSharedMemStUniformCudaHipBuiltIn const&) -> void
50  {
51  // Nothing to do. CUDA/HIP block shared memory is automatically freed when all threads left the block.
52  }
53  };
54  } // namespace trait
55 
56 # endif
57 
58 } // namespace alpaka
59 
60 #endif
The alpaka accelerator library.
Tag used in class inheritance hierarchies that describes that a specific concept (TConcept) is implem...
Definition: Concepts.hpp:15
static __device__ auto declareVar(BlockSharedMemStUniformCudaHipBuiltIn const &) -> T &
The block shared static memory variable allocation operation trait.
Definition: Traits.hpp:23
static __device__ auto freeVars(BlockSharedMemStUniformCudaHipBuiltIn const &) -> void
The block shared static memory free operation trait.
Definition: Traits.hpp:26