Skip to content

Commit 57bdbe3

Browse files
0x12CC0x12CC
authored
[SYCL] Implement sycl_khr_{static, dynamic}_addrspace_cast extensions (#18521)
Implement the extensions defined in KhronosGroup/SYCL-Docs#650. --------- Signed-off-by: Michael Aziz <[email protected]>
1 parent d3bc93a commit 57bdbe3

File tree

6 files changed

+269
-1
lines changed

6 files changed

+269
-1
lines changed

‎sycl/include/sycl/khr/dynamic_addrspace_cast.hpp

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
//==-- dynamic_addrspace_cast.hpp --- KHR dynamic addrspace cast 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+
10+
#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
11+
12+
#include <sycl/access/access.hpp>
13+
#include <sycl/ext/oneapi/experimental/address_cast.hpp>
14+
#include <sycl/multi_ptr.hpp>
15+
16+
namespace sycl {
17+
inline namespace _V1 {
18+
namespace khr {
19+
20+
template <access::address_space Space, typename ElementType>
21+
multi_ptr<ElementType, Space, access::decorated::no>
22+
dynamic_addrspace_cast(ElementType *ptr) {
23+
return ext::oneapi::experimental::dynamic_address_cast<Space>(ptr);
24+
}
25+
26+
template <access::address_space Space, typename ElementType,
27+
access::decorated DecorateAddress>
28+
multi_ptr<ElementType, Space, DecorateAddress> dynamic_addrspace_cast(
29+
multi_ptr<ElementType, access::address_space::generic_space,
30+
DecorateAddress>
31+
ptr) {
32+
return ext::oneapi::experimental::dynamic_address_cast<Space>(ptr);
33+
}
34+
35+
} // namespace khr
36+
} // namespace _V1
37+
} // namespace sycl
38+
39+
#endif

‎sycl/include/sycl/khr/static_addrspace_cast.hpp

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
//==--- static_addrspace_cast.hpp --- KHR static addrspace cast 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+
10+
#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
11+
12+
#include <sycl/access/access.hpp>
13+
#include <sycl/ext/oneapi/experimental/address_cast.hpp>
14+
#include <sycl/multi_ptr.hpp>
15+
16+
namespace sycl {
17+
inline namespace _V1 {
18+
namespace khr {
19+
20+
template <access::address_space Space, typename ElementType>
21+
multi_ptr<ElementType, Space, access::decorated::no>
22+
static_addrspace_cast(ElementType *ptr) {
23+
return ext::oneapi::experimental::static_address_cast<Space>(ptr);
24+
}
25+
26+
template <access::address_space Space, typename ElementType,
27+
access::decorated DecorateAddress>
28+
multi_ptr<ElementType, Space, DecorateAddress> static_addrspace_cast(
29+
multi_ptr<ElementType, access::address_space::generic_space,
30+
DecorateAddress>
31+
ptr) {
32+
return ext::oneapi::experimental::static_address_cast<Space>(ptr);
33+
}
34+
35+
} // namespace khr
36+
} // namespace _V1
37+
} // namespace sycl
38+
39+
#endif

‎sycl/include/sycl/sycl.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -122,6 +122,8 @@
122122
#include <sycl/ext/oneapi/virtual_mem/physical_mem.hpp>
123123
#include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
124124
#include <sycl/ext/oneapi/weak_object.hpp>
125+
#include <sycl/khr/dynamic_addrspace_cast.hpp>
125126
#include <sycl/khr/free_function_commands.hpp>
126127
#include <sycl/khr/group_interface.hpp>
128+
#include <sycl/khr/static_addrspace_cast.hpp>
127129
#include <sycl/khr/work_item_queries.hpp>

‎sycl/source/feature_test.hpp.in

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -112,7 +112,7 @@ inline namespace _V1 {
112112
#define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1
113113
#define SYCL_EXT_ONEAPI_WORK_GROUP_MEMORY 1
114114
#define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1
115-
#define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1
115+
#define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1
116116
#define SYCL_EXT_ONEAPI_NUM_COMPUTE_UNITS 1
117117
#define SYCL_EXT_ONEAPI_DEVICE_IMAGE_BACKEND_CONTENT 1
118118
#define SYCL_EXT_ONEAPI_CURRENT_DEVICE 1
@@ -129,6 +129,13 @@ inline namespace _V1 {
129129
#define SYCL_KHR_WORK_ITEM_QUERIES 1
130130
#endif
131131

132+
// Unfinished KHR extensions. These extensions are only available if the
133+
// __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS macro is defined.
134+
#ifdef __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
135+
#define SYCL_KHR_STATIC_ADDRSPACE_CAST 1
136+
#define SYCL_KHR_DYNAMIC_ADDRSPACE_CAST 1
137+
#endif
138+
132139
#ifndef __has_include
133140
#define __has_include(x) 0
134141
#endif

‎sycl/test-e2e/AddressCast/khr_dynamic_addrspace_cast.cpp

Lines changed: 109 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,109 @@
1+
//==------ khr_dynamic_addrspace_cast.cpp - dynamic addrspace cast test ----==//
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+
9+
// Issue with OpenCL CPU runtime implementation of OpGenericCastToPtrExplicit
10+
// FPGA emulator affected by same issue as OpenCL CPU runtime
11+
// UNSUPPORTED: cpu, accelerator
12+
// UNSUPPORTED-TRACKER: CMPLRLLVM-37365
13+
// RUN: %{build} -o %t.out
14+
// RUN: %{run} %t.out
15+
16+
// Depends on SPIR-V Backend & run-time drivers version.
17+
// XFAIL: spirv-backend && gpu
18+
// XFAIL-TRACKER: CMPLRLLVM-64705
19+
20+
#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
21+
22+
#include <sycl/detail/core.hpp>
23+
#include <sycl/khr/dynamic_addrspace_cast.hpp>
24+
25+
int main() {
26+
27+
sycl::queue Queue;
28+
29+
sycl::range<1> NItems{1};
30+
31+
sycl::buffer<int, 1> GlobalBuffer{NItems};
32+
sycl::buffer<bool, 1> ResultBuffer{NItems};
33+
34+
Queue
35+
.submit([&](sycl::handler &cgh) {
36+
auto GlobalAccessor =
37+
GlobalBuffer.get_access<sycl::access::mode::read_write>(cgh);
38+
auto LocalAccessor = sycl::local_accessor<int>(1, cgh);
39+
auto ResultAccessor =
40+
ResultBuffer.get_access<sycl::access::mode::write>(cgh);
41+
cgh.parallel_for<class Kernel>(
42+
sycl::nd_range<1>(NItems, 1), [=](sycl::nd_item<1> Item) {
43+
bool Success = true;
44+
size_t Index = Item.get_global_id(0);
45+
46+
int *RawGlobalPointer = &GlobalAccessor[Index];
47+
{
48+
auto GlobalPointer = sycl::khr::dynamic_addrspace_cast<
49+
sycl::access::address_space::global_space>(
50+
RawGlobalPointer);
51+
auto LocalPointer = sycl::khr::dynamic_addrspace_cast<
52+
sycl::access::address_space::local_space>(RawGlobalPointer);
53+
auto PrivatePointer = sycl::khr::dynamic_addrspace_cast<
54+
sycl::access::address_space::private_space>(
55+
RawGlobalPointer);
56+
Success &= reinterpret_cast<size_t>(RawGlobalPointer) ==
57+
reinterpret_cast<size_t>(GlobalPointer.get_raw());
58+
Success &= LocalPointer.get_raw() == nullptr;
59+
Success &= PrivatePointer.get_raw() == nullptr;
60+
}
61+
62+
int *RawLocalPointer = &LocalAccessor[0];
63+
{
64+
auto GlobalPointer = sycl::khr::dynamic_addrspace_cast<
65+
sycl::access::address_space::global_space>(RawLocalPointer);
66+
auto LocalPointer = sycl::khr::dynamic_addrspace_cast<
67+
sycl::access::address_space::local_space>(RawLocalPointer);
68+
auto PrivatePointer = sycl::khr::dynamic_addrspace_cast<
69+
sycl::access::address_space::private_space>(
70+
RawLocalPointer);
71+
Success &= GlobalPointer.get_raw() == nullptr;
72+
Success &= reinterpret_cast<size_t>(RawLocalPointer) ==
73+
reinterpret_cast<size_t>(LocalPointer.get_raw());
74+
Success &= PrivatePointer.get_raw() == nullptr;
75+
}
76+
77+
int PrivateVariable = 0;
78+
int *RawPrivatePointer = &PrivateVariable;
79+
{
80+
auto GlobalPointer = sycl::khr::dynamic_addrspace_cast<
81+
sycl::access::address_space::global_space>(
82+
RawPrivatePointer);
83+
auto LocalPointer = sycl::khr::dynamic_addrspace_cast<
84+
sycl::access::address_space::local_space>(
85+
RawPrivatePointer);
86+
auto PrivatePointer = sycl::khr::dynamic_addrspace_cast<
87+
sycl::access::address_space::private_space>(
88+
RawPrivatePointer);
89+
Success &= GlobalPointer.get_raw() == nullptr;
90+
Success &= LocalPointer.get_raw() == nullptr;
91+
Success &= reinterpret_cast<size_t>(RawPrivatePointer) ==
92+
reinterpret_cast<size_t>(PrivatePointer.get_raw());
93+
}
94+
95+
ResultAccessor[Index] = Success;
96+
});
97+
})
98+
.wait();
99+
100+
bool Success = true;
101+
{
102+
auto ResultAccessor = sycl::host_accessor(ResultBuffer);
103+
for (int i = 0; i < NItems.size(); ++i) {
104+
Success &= ResultAccessor[i];
105+
};
106+
}
107+
108+
return (Success) ? 0 : -1;
109+
}

‎sycl/test-e2e/AddressCast/khr_static_addrspace_cast.cpp

Lines changed: 72 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,72 @@
1+
//==-------- khr_static_addrspace_cast.cpp - static addrspace cast test ----==//
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+
9+
// RUN: %{build} -o %t.out
10+
// RUN: %{run} %t.out
11+
12+
#define __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS
13+
14+
#include <sycl/detail/core.hpp>
15+
#include <sycl/khr/static_addrspace_cast.hpp>
16+
17+
int main() {
18+
19+
sycl::queue Queue;
20+
21+
sycl::range<1> NItems{1};
22+
23+
sycl::buffer<int, 1> GlobalBuffer{NItems};
24+
sycl::buffer<bool, 1> ResultBuffer{NItems};
25+
26+
Queue
27+
.submit([&](sycl::handler &cgh) {
28+
auto GlobalAccessor =
29+
GlobalBuffer.get_access<sycl::access::mode::read_write>(cgh);
30+
auto LocalAccessor = sycl::local_accessor<int>(1, cgh);
31+
auto ResultAccessor =
32+
ResultBuffer.get_access<sycl::access::mode::write>(cgh);
33+
cgh.parallel_for<class Kernel>(
34+
sycl::nd_range<1>(NItems, 1), [=](sycl::nd_item<1> Item) {
35+
bool Success = true;
36+
size_t Index = Item.get_global_id(0);
37+
38+
int *RawGlobalPointer = &GlobalAccessor[Index];
39+
auto GlobalPointer = sycl::khr::static_addrspace_cast<
40+
sycl::access::address_space::global_space>(RawGlobalPointer);
41+
Success &= reinterpret_cast<size_t>(RawGlobalPointer) ==
42+
reinterpret_cast<size_t>(GlobalPointer.get_raw());
43+
44+
int *RawLocalPointer = &LocalAccessor[0];
45+
auto LocalPointer = sycl::khr::static_addrspace_cast<
46+
sycl::access::address_space::local_space>(RawLocalPointer);
47+
Success &= reinterpret_cast<size_t>(RawLocalPointer) ==
48+
reinterpret_cast<size_t>(LocalPointer.get_raw());
49+
50+
int PrivateVariable = 0;
51+
int *RawPrivatePointer = &PrivateVariable;
52+
auto PrivatePointer = sycl::khr::static_addrspace_cast<
53+
sycl::access::address_space::private_space>(
54+
RawPrivatePointer);
55+
Success &= reinterpret_cast<size_t>(RawPrivatePointer) ==
56+
reinterpret_cast<size_t>(PrivatePointer.get_raw());
57+
58+
ResultAccessor[Index] = Success;
59+
});
60+
})
61+
.wait();
62+
63+
bool Success = true;
64+
{
65+
auto ResultAccessor = sycl::host_accessor(ResultBuffer);
66+
for (int i = 0; i < NItems.size(); ++i) {
67+
Success &= ResultAccessor[i];
68+
};
69+
}
70+
71+
return (Success) ? 0 : -1;
72+
}

0 commit comments

Comments
 (0)