|
14 | 14 | #include <CL/sycl/detail/common.hpp>
|
15 | 15 | #include <CL/sycl/detail/generic_type_traits.hpp>
|
16 | 16 | #include <CL/sycl/detail/helpers.hpp>
|
| 17 | +#include <CL/sycl/detail/spirv.hpp> |
17 | 18 | #include <CL/sycl/device_event.hpp>
|
18 | 19 | #include <CL/sycl/h_item.hpp>
|
19 | 20 | #include <CL/sycl/id.hpp>
|
20 | 21 | #include <CL/sycl/memory_enums.hpp>
|
21 | 22 | #include <CL/sycl/pointers.hpp>
|
22 | 23 | #include <CL/sycl/range.hpp>
|
23 |
| -#include <CL/sycl/detail/spirv.hpp> |
24 | 24 | #include <stdexcept>
|
25 | 25 | #include <type_traits>
|
26 | 26 |
|
@@ -51,7 +51,8 @@ static inline void workGroupBarrier(memory_scope FenceScope) {
|
51 | 51 | __spv::MemorySemanticsMask::WorkgroupMemory |
|
52 | 52 | __spv::MemorySemanticsMask::CrossWorkgroupMemory);
|
53 | 53 | #else
|
54 |
| - throw sycl::runtime_error("Barriers are not supported on host device", PI_INVALID_DEVICE); |
| 54 | + throw sycl::runtime_error("Barriers are not supported on host device", |
| 55 | + PI_INVALID_DEVICE); |
55 | 56 | #endif
|
56 | 57 | }
|
57 | 58 |
|
@@ -450,13 +451,16 @@ template <int Dims> group<Dims> this_group() {
|
450 | 451 | template <typename Group>
|
451 | 452 | void group_barrier(Group G, memory_scope FenceScope = Group::fence_scope);
|
452 | 453 |
|
453 |
| -template <> inline void group_barrier<group<1>>(group<1>, memory_scope FenceScope) { |
| 454 | +template <> |
| 455 | +inline void group_barrier<group<1>>(group<1>, memory_scope FenceScope) { |
454 | 456 | detail::workGroupBarrier(FenceScope);
|
455 | 457 | }
|
456 |
| -template <> inline void group_barrier<group<2>>(group<2>, memory_scope FenceScope) { |
| 458 | +template <> |
| 459 | +inline void group_barrier<group<2>>(group<2>, memory_scope FenceScope) { |
457 | 460 | detail::workGroupBarrier(FenceScope);
|
458 | 461 | }
|
459 |
| -template <> inline void group_barrier<group<3>>(group<3>, memory_scope FenceScope) { |
| 462 | +template <> |
| 463 | +inline void group_barrier<group<3>>(group<3>, memory_scope FenceScope) { |
460 | 464 | detail::workGroupBarrier(FenceScope);
|
461 | 465 | }
|
462 | 466 |
|
|
0 commit comments