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 
14 #include <sycl/exception.hpp>
15 #include <sycl/group.hpp>
16 #include <sycl/multi_ptr.hpp>
17 
18 #include <cstdint>
19 #include <type_traits>
20 #include <utility>
21 
22 namespace sycl {
24 namespace ext::oneapi {
25 template <typename T, typename Group>
27  std::is_trivially_destructible<T>::value &&
28  sycl::detail::is_group<Group>::value,
29  multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
31  (void)g;
32 #ifdef __SYCL_DEVICE_ONLY__
33  __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
34  __sycl_allocateLocalMemory(sizeof(T), alignof(T));
35  return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
36 #else
38  "sycl_ext_oneapi_local_memory extension is not supported on host device",
39  PI_ERROR_INVALID_OPERATION);
40 #endif
41 }
42 
43 template <typename T, typename Group, typename... Args>
45  std::is_trivially_destructible<T>::value &&
46  sycl::detail::is_group<Group>::value,
48  __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args) {
49  (void)g;
50 #ifdef __SYCL_DEVICE_ONLY__
51  __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
52  __sycl_allocateLocalMemory(sizeof(T), alignof(T));
53 
54  // TODO switch to using group::get_local_linear_id here once it's implemented
55  id<3> Id = __spirv::initLocalInvocationId<3, id<3>>();
56  if (Id == id<3>(0, 0, 0))
57  new (AllocatedMem) T(std::forward<Args>(args)...);
59  return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
60 #else
61  // Silence unused variable warning
62  [&args...] {}();
64  "sycl_ext_oneapi_local_memory extension is not supported on host device",
65  PI_ERROR_INVALID_OPERATION);
66 #endif
67 }
68 } // namespace ext::oneapi
69 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
70 } // namespace sycl
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: multi_ptr.hpp:78
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL_ALWAYS_INLINE
typename std::enable_if< B, T >::type enable_if_t
static void workGroupBarrier()
Definition: group.hpp:33
__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< T >::value &&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< T >::value &&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)
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14