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 <CL/sycl/exception.hpp>
15 #include <CL/sycl/group.hpp>
16 #include <CL/sycl/multi_ptr.hpp>
17 
18 #include <cstdint>
19 #include <type_traits>
20 #include <utility>
21 
23 namespace sycl {
24 namespace ext {
25 namespace oneapi {
26 template <typename T, typename Group>
27 std::enable_if_t<std::is_trivially_destructible<T>::value &&
28  sycl::detail::is_group<Group>::value,
29  multi_ptr<T, access::address_space::local_space>>
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",
40 #endif
41 }
42 
43 template <typename T, typename Group, typename... Args>
44 std::enable_if_t<std::is_trivially_destructible<T>::value &&
47  __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&... args) {
48  (void)g;
49 #ifdef __SYCL_DEVICE_ONLY__
50  __attribute__((opencl_local)) std::uint8_t *AllocatedMem =
51  __sycl_allocateLocalMemory(sizeof(T), alignof(T));
52 
53  // TODO switch to using group::get_local_linear_id here once it's implemented
54  id<3> Id = __spirv::initLocalInvocationId<3, id<3>>();
55  if (Id == id<3>(0, 0, 0))
56  new (AllocatedMem) T(std::forward<Args>(args)...);
58  return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
59 #else
60  // Silence unused variable warning
61  [&args...] {}();
63  "sycl_ext_oneapi_local_memory extension is not supported on host device",
65 #endif
66 }
67 } // namespace oneapi
68 } // namespace ext
69 } // namespace sycl
70 } // __SYCL_INLINE_NAMESPACE(cl)
T
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:88
type_traits.hpp
defines_elementary.hpp
cl::sycl::id< 3 >
cl::sycl::ext::oneapi::group_local_memory
std::enable_if_t< std::is_trivially_destructible< T >::value &&sycl::detail::is_group< Group >::value, multi_ptr< T, access::address_space::local_space > > __SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&... args)
Definition: group_local_memory.hpp:47
cl::sycl::detail::workGroupBarrier
static void workGroupBarrier()
Definition: group.hpp:33
cl::sycl::errc::feature_not_supported
@ feature_not_supported
spirv_vars.hpp
sycl
Definition: invoke_simd.hpp:68
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:32
cl::sycl::ext::oneapi::group_local_memory_for_overwrite
std::enable_if_t< std::is_trivially_destructible< T >::value &&sycl::detail::is_group< Group >::value, multi_ptr< T, access::address_space::local_space > > __SYCL_ALWAYS_INLINE group_local_memory_for_overwrite(Group g)
Definition: group_local_memory.hpp:30
group.hpp
__SYCL_ALWAYS_INLINE
#define __SYCL_ALWAYS_INLINE
Definition: defines_elementary.hpp:29
multi_ptr.hpp
sycl_fe_intrins.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::is_group
Definition: type_traits.hpp:31
sycl::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:293
exception.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12