[SYCL] Implement work group memory extension #15178
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,103 @@ | ||
| //===-------------------- work_group_memory.hpp ---------------------------===// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #pragma once | ||
|
|
||
| #include <type_traits> | ||
|
|
||
| namespace sycl { | ||
| inline namespace _V1 { | ||
| namespace detail { | ||
| template <typename T> struct is_unbounded_array : std::false_type {}; | ||
|
|
||
| template <typename T> struct is_unbounded_array<T[]> : std::true_type {}; | ||
|
|
||
| template <typename T> | ||
| inline constexpr bool is_unbounded_array_v = is_unbounded_array<T>::value; | ||
|
|
||
| class work_group_memory_impl { | ||
| public: | ||
| work_group_memory_impl() : wgm_size{0}, buffer_size{0} {} | ||
| work_group_memory_impl(const work_group_memory_impl &rhs) = default; | ||
| work_group_memory_impl & | ||
| operator=(const work_group_memory_impl &rhs) = default; | ||
| work_group_memory_impl(size_t wgm_size, size_t buffer_size) | ||
| : wgm_size{wgm_size}, buffer_size{buffer_size} {} | ||
| size_t wgm_size; | ||
| size_t buffer_size; | ||
|
There was a problem hiding this comment. I haven't looked at this closely, but this looks suspicious. The It looks like There was a problem hiding this comment. Initial testing shows that the L0 kernel only ends up with one parameter, but I will make this into a proper test. There was a problem hiding this comment. I wonder why there is only one Level Zero parameter. Is it because When I look at the code, I don't see any uses of If all we need is Or, you could use There was a problem hiding this comment. So, my understanding is that the There was a problem hiding this comment. The relevant lines are 793-798 in There was a problem hiding this comment. OK, thanks for the explanation! There was a problem hiding this comment. @lbushi25 is right that this will only result in a single kernel argument, which is due to it being a "SYCL special class" so the arguments of its All that said however, I also have some concerns here:
#ifdef __SYCL_DEVICE_ONLY__
decoratedPtr ptr;
// To ensure we have the same object size on host and device, we add padding.
[[maybe_unused]] char padding[sizeof(std::shared_ptr<work_group_memory_impl>) - sizeof(decoratedPtr)];
#else
std::shared_ptr<work_group_memory_impl> impl;
#endifNote that if we're sure we only ever need the two There was a problem hiding this comment.
But I was also confused at first, so I think it worth explicitly declaring members of the base class as |
||
| }; | ||
|
|
||
| inline size_t getWorkGroupMemoryOwnSize(detail::work_group_memory_impl *wgm) { | ||
| return wgm->wgm_size; | ||
| } | ||
|
|
||
| // The following 3 functions help us get the address of the first element of a | ||
| // multi-dimensional array, be it bounded or unbounded. A scalar is also | ||
| // included. In that case, it just returns the address of the scalar. | ||
| template <typename DataT> auto getData(DataT &scalar) { return &scalar; } | ||
|
|
||
| template <typename DataT, size_t N> auto getData(DataT (&bounded_arr)[N]) { | ||
| return getData(bounded_arr[0]); | ||
| } | ||
|
|
||
| template <typename DataT> auto getData(DataT (&unbounded_arr)[]) { | ||
| return getData(unbounded_arr[0]); | ||
| } | ||
|
|
||
| } // namespace detail | ||
|
|
||
| namespace ext::oneapi::experimental { | ||
| template <typename DataT, typename PropertyListT = empty_properties_t> | ||
| class __SYCL_SPECIAL_CLASS __SYCL_TYPE(work_group_memory) work_group_memory | ||
| : sycl::detail::work_group_memory_impl { | ||
| public: | ||
| using value_type = std::remove_all_extents_t<DataT>; | ||
|
|
||
| private: | ||
| using decoratedPtr = typename sycl::detail::DecoratedType< | ||
| DataT, access::address_space::local_space>::type *; | ||
|
||
|
|
||
| public: | ||
| work_group_memory() = default; | ||
| work_group_memory(const work_group_memory &rhs) = default; | ||
| work_group_memory &operator=(const work_group_memory &rhs) = default; | ||
| template <typename T = DataT, | ||
| typename = std::enable_if_t<!sycl::detail::is_unbounded_array_v<T>>> | ||
| work_group_memory(handler &) | ||
| : sycl::detail::work_group_memory_impl(sizeof(work_group_memory), | ||
| sizeof(DataT)) {} | ||
| template <typename T = DataT, | ||
| typename = std::enable_if_t<sycl::detail::is_unbounded_array_v<T>>> | ||
| work_group_memory(size_t num, handler &) | ||
| : sycl::detail::work_group_memory_impl( | ||
| sizeof(work_group_memory), | ||
| num * sizeof(std::remove_extent_t<DataT>)) {} | ||
| template <access::decorated IsDecorated = access::decorated::no> | ||
| multi_ptr<value_type, access::address_space::local_space, IsDecorated> | ||
| get_multi_ptr() const { | ||
| return sycl::address_space_cast<access::address_space::local_space, | ||
| IsDecorated, value_type>( | ||
| sycl::detail::getData(*ptr)); | ||
|
||
| } | ||
| DataT *operator&() const { return ptr; } | ||
| operator DataT &() const { return *(this->operator&()); } | ||
|
||
| template <typename T = DataT, | ||
| typename = std::enable_if_t<!std::is_array_v<T>>> | ||
| const work_group_memory &operator=(const DataT &value) const { | ||
| *ptr = value; | ||
| return *this; | ||
| } | ||
| #ifdef __SYCL_DEVICE_ONLY__ | ||
| void __init(decoratedPtr ptr) { this->ptr = ptr; } | ||
| #endif | ||
| private: | ||
| decoratedPtr ptr; | ||
| }; | ||
| } // namespace ext::oneapi::experimental | ||
| } // namespace _V1 | ||
| } // namespace sycl | ||
Uh oh!
There was an error while loading. Please reload this page.