Skip to content

Commit b50440e

Browse files
authored
[SYCL] Use aggregate initialization for group_local_memory arrays (#10367)
Signed-off-by: Michael Aziz <[email protected]>
1 parent aaa2e13 commit b50440e

File tree

2 files changed

+24
-1
lines changed

2 files changed

+24
-1
lines changed

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,7 @@ std::enable_if_t<
5454
__attribute__((opencl_local)) std::uint8_t *AllocatedMem =
5555
__sycl_allocateLocalMemory(sizeof(T), alignof(T));
5656
if (g.get_local_linear_id() == 0)
57-
new (AllocatedMem) T(std::forward<Args>(args)...);
57+
new (AllocatedMem) T{std::forward<Args>(args)...};
5858
sycl::detail::workGroupBarrier();
5959
return reinterpret_cast<__attribute__((opencl_local)) T *>(AllocatedMem);
6060
#else
Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
4+
#include <CL/sycl.hpp>
5+
6+
constexpr int N = 5;
7+
8+
int main() {
9+
sycl::queue q;
10+
int *ptr = sycl::malloc_shared<int>(N, q);
11+
q.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1> it) {
12+
auto g = it.get_group();
13+
auto mem = sycl::ext::oneapi::group_local_memory<int[N]>(g, 1, 2, 3, 4, 5);
14+
auto ref = *mem;
15+
for (int i = 0; i < N; ++i) {
16+
ptr[i] = ref[i];
17+
}
18+
}).wait();
19+
for (int i = 0; i < N; ++i) {
20+
assert(ptr[i] == (i + 1));
21+
}
22+
sycl::free(ptr, q);
23+
}

0 commit comments

Comments
 (0)