Skip to content

Commit 442fe98

Browse files
committed
add back __SYCL_ALWAYS_INLINE
1 parent db2ad4a commit 442fe98

File tree

2 files changed

+3
-3
lines changed

2 files changed

+3
-3
lines changed

sycl/include/sycl/ext/oneapi/group_local_memory.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ template <typename T, typename Group>
2727
std::enable_if_t<
2828
std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
2929
multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
30-
group_local_memory_for_overwrite(Group g) {
30+
__SYCL_ALWAYS_INLINE group_local_memory_for_overwrite(Group g) {
3131
(void)g;
3232
#ifdef __SYCL_DEVICE_ONLY__
3333
__attribute__((opencl_local)) std::uint8_t *AllocatedMem =
@@ -53,7 +53,7 @@ template <typename T, typename Group, typename... Args>
5353
std::enable_if_t<
5454
std::is_trivially_destructible_v<T> && sycl::detail::is_group<Group>::value,
5555
multi_ptr<T, access::address_space::local_space, access::decorated::legacy>>
56-
group_local_memory(Group g, Args &&...args) {
56+
__SYCL_ALWAYS_INLINE group_local_memory(Group g, Args &&...args) {
5757
#ifdef __SYCL_DEVICE_ONLY__
5858
__attribute__((opencl_local)) std::uint8_t *AllocatedMem =
5959
__sycl_allocateLocalMemory(sizeof(T), alignof(T));

sycl/include/syclcompat/memory.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,7 +72,7 @@ template <typename AllocT>
7272
#ifdef __SYCL_DEVICE_ONLY__
7373
[[__sycl_detail__::add_ir_attributes_function("sycl-forceinline", true)]]
7474
#endif
75-
auto *local_mem() {
75+
__SYCL_ALWAYS_INLINE auto *local_mem() {
7676
sycl::multi_ptr<AllocT, sycl::access::address_space::local_space>
7777
As_multi_ptr =
7878
sycl::ext::oneapi::group_local_memory_for_overwrite<AllocT>(

0 commit comments

Comments
 (0)