DPC++ Runtime
Runtime libraries for oneAPI DPC++
group_local_memory.hpp
Go to the documentation of this file.
1 //==----- group_local_memory.hpp --- SYCL group local memory extension -----==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 #pragma once
9 
10 #include <sycl/access/access.hpp> // for address_space, decorated
11 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
12 #include <sycl/detail/pi.h> // for PI_ERROR_INVALID_OPERA...
13 #include <sycl/detail/type_traits.hpp> // for is_group
14 #include <sycl/exception.hpp> // for exception
15 #include <sycl/ext/intel/usm_pointers.hpp> // for multi_ptr
16 
17 #include <type_traits> // for enable_if_t
18 
19 #ifdef __SYCL_DEVICE_ONLY__
20 // Request a fixed-size allocation in local address space at kernel scope.
21 extern "C" __DPCPP_SYCL_EXTERNAL __attribute__((opencl_local)) std::uint8_t *
22 __sycl_allocateLocalMemory(std::size_t Size, std::size_t Alignment);
23 #endif
24 
25 namespace sycl {
26 inline namespace _V1 {
27 namespace ext::oneapi {
28 template <typename T, typename Group>
29 std::enable_if_t<
30  std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
31  multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
33  (void)g;
34 #ifdef __SYCL_DEVICE_ONLY__
35  __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
36  __sycl_allocateLocalMemory(sizeof(T), alignof(T));
37  // If the type is non-trivial we need to default initialize it.
38  if constexpr (!std::is_trivial_v<T>) {
39  if (g.get_local_linear_id() == 0)
40  new (AllocatedMem) T; // Default initialize.
42  }
43  return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
44 #else
45  throw sycl::exception(
46  sycl::errc::feature_not_supported,
47  "sycl_ext_oneapi_local_memory extension is not supported on host");
48 #endif
49 }
50 
51 template <typename T, typename Group, typename... Args>
52 std::enable_if_t<
53  std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
55  __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args) {
56 #ifdef __SYCL_DEVICE_ONLY__
57  __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
58  __sycl_allocateLocalMemory(sizeof(T), alignof(T));
59  if (g.get_local_linear_id() == 0)
60  new (AllocatedMem) T{std::forward<Args>(args)...};
62  return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
63 #else
64  // Silence unused variable warning
65  (void)g;
66  [&args...] {}();
67  throw sycl::exception(
68  sycl::errc::feature_not_supported,
69  "sycl_ext_oneapi_local_memory extension is not supported on host");
70 #endif
71 }
72 } // namespace ext::oneapi
73 } // namespace _V1
74 } // namespace sycl
#define __SYCL_ALWAYS_INLINE
#define __DPCPP_SYCL_EXTERNAL
void workGroupBarrier()
Definition: group.hpp:42
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
std::enable_if_t< std::is_trivially_destructible_v< T > &&sycl::detail::is_group< Group >::value, multi_ptr< T, access::address_space::local_space, access::decorated::legacy > > __SYCL_ALWAYS_INLINE group_local_memory_for_overwrite(Group g)
std::enable_if_t< std::is_trivially_destructible_v< T > &&sycl::detail::is_group< Group >::value, multi_ptr< T, access::address_space::local_space, access::decorated::legacy > > __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args)
Definition: access.hpp:18