Skip to content

Commit 0b8cd4e

Browse files
JackAKirkJackAKirk
committed
Barebones impl for testing new prop.
Signed-off-by: JackAKirk <[email protected]>
1 parent b0d9167 commit 0b8cd4e

File tree

7 files changed

+152
-69
lines changed

7 files changed

+152
-69
lines changed

‎clang/include/clang/Basic/Attr.td‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1805,6 +1805,15 @@ def SYCLDeviceGlobal: InheritableAttr {
18051805
let SimpleHandler = 1;
18061806
}
18071807

1808+
def SYCLDeviceConstant: InheritableAttr {
1809+
let Spellings = [CXX11<"__sycl_detail__", "device_constant">];
1810+
let Subjects = SubjectList<[CXXRecord], ErrorDiag>;
1811+
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
1812+
// Only used internally by SYCL implementation
1813+
let Documentation = [Undocumented];
1814+
let SimpleHandler = 1;
1815+
}
1816+
18081817
def SYCLGlobalVariableAllowed : InheritableAttr {
18091818
let Spellings = [CXX11<"__sycl_detail__", "global_variable_allowed">];
18101819
let Subjects = SubjectList<[CXXRecord], ErrorDiag>;

‎clang/lib/CodeGen/CodeGenModule.cpp‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5774,10 +5774,7 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
57745774

57755775
if (getTriple().isNVPTX() || getTriple().isAMDGPU()) {
57765776
const RecordDecl *RD = D->getType()->getAsRecordDecl();
5777-
if (RD && RD->hasAttr<SYCLDeviceGlobalAttr>()) {
5778-
clang::CXXRecordDecl *record = D->getType()->getAsCXXRecordDecl();
5779-
const auto *Spec = cast<ClassTemplateSpecializationDecl>(record);
5780-
if (Spec->getTemplateArgs().get(0).getAsType().isConstQualified())
5777+
if (RD && RD->hasAttr<SYCLDeviceConstantAttr>()) {
57815778
return LangAS::opencl_constant;
57825779
}
57835780
}

‎sycl/include/sycl/ext/oneapi/device_global/device_global.hpp‎

Lines changed: 98 additions & 62 deletions
Original file line numberDiff line numberDiff line change
@@ -192,7 +192,7 @@ class device_global_base<
192192

193193
} // namespace detail
194194

195-
template <typename T, typename PropertyListT = empty_properties_t>
195+
template <typename T, typename PropertyListT, typename Condition>
196196
class
197197
#ifdef __SYCL_DEVICE_ONLY__
198198
// FIXME: Temporary work-around. Remove when fixed.
@@ -204,20 +204,82 @@ class
204204
"Property list is invalid.");
205205
};
206206

207+
#define DEVICE_GLOBAL_COMMON_METHODS() \
208+
using detail::device_global_base< \
209+
T, detail::properties_t<Props...>>::device_global_base; \
210+
\
211+
device_global(const device_global &) = delete; \
212+
device_global(const device_global &&) = delete; \
213+
device_global &operator=(const device_global &) = delete; \
214+
device_global &operator=(const device_global &&) = delete; \
215+
const T &get() const noexcept { \
216+
__SYCL_HOST_NOT_SUPPORTED("get()") \
217+
return *this->get_ptr(); \
218+
} \
219+
\
220+
operator const T &() const noexcept { \
221+
__SYCL_HOST_NOT_SUPPORTED("Implicit conversion of device_global to T") \
222+
return get(); \
223+
} \
224+
\
225+
template <class RelayT = T> \
226+
std::remove_reference_t< \
227+
decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> \
228+
&operator[](std::ptrdiff_t idx) noexcept { \
229+
__SYCL_HOST_NOT_SUPPORTED("Subscript operator") \
230+
return (*this->get_ptr())[idx]; \
231+
} \
232+
\
233+
template <class RelayT = T> \
234+
const std::remove_reference_t< \
235+
decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> \
236+
&operator[](std::ptrdiff_t idx) const noexcept { \
237+
__SYCL_HOST_NOT_SUPPORTED("Subscript operator") \
238+
return (*this->get_ptr())[idx]; \
239+
} \
240+
\
241+
template <class RelayT = T> \
242+
std::enable_if_t<detail::HasArrowOperator<RelayT>::value || \
243+
std::is_pointer_v<RelayT>, \
244+
RelayT> \
245+
&operator->() noexcept { \
246+
__SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global") \
247+
return *this->get_ptr(); \
248+
} \
249+
\
250+
template <class RelayT = T> \
251+
std::enable_if_t<detail::HasArrowOperator<RelayT>::value || \
252+
std::is_pointer_v<RelayT>, \
253+
const RelayT> \
254+
&operator->() const noexcept { \
255+
__SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global") \
256+
return *this->get_ptr(); \
257+
} \
258+
\
259+
template <typename propertyT> static constexpr bool has_property() { \
260+
return property_list_t::template has_property<propertyT>(); \
261+
} \
262+
\
263+
template <typename propertyT> static constexpr auto get_property() { \
264+
return property_list_t::template get_property<propertyT>(); \
265+
}
266+
207267
template <typename T, typename... Props>
208268
class
209269
#ifdef __SYCL_DEVICE_ONLY__
210270
[[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global,
271+
__sycl_detail__::device_constant,
211272
__sycl_detail__::add_ir_attributes_global_variable(
212273
"sycl-device-global-size",
213274
__SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::name..., sizeof(T),
214275
__SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::value...)]]
215276
#endif
216-
device_global<T, detail::properties_t<Props...>>
277+
device_global<T, detail::properties_t<Props...>,
278+
std::enable_if_t<detail::properties_t<
279+
Props...>::template has_property<device_constant_key>()>>
217280
: public detail::device_global_base<T, detail::properties_t<Props...>> {
218281

219282
using property_list_t = detail::properties_t<Props...>;
220-
using base_t = detail::device_global_base<T, property_list_t>;
221283

222284
public:
223285
using element_type = std::remove_extent_t<T>;
@@ -233,22 +295,43 @@ class
233295
static_assert(is_property_list<property_list_t>::value,
234296
"Property list is invalid.");
235297

236-
// Inherit the base class' constructors
237-
using detail::device_global_base<T, property_list_t>::device_global_base;
298+
DEVICE_GLOBAL_COMMON_METHODS()
299+
};
238300

239-
constexpr device_global(const device_global &DG)
240-
: base_t(static_cast<const base_t &>(DG)) {}
301+
template <typename T, typename... Props>
302+
class
303+
#ifdef __SYCL_DEVICE_ONLY__
304+
[[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global,
305+
__sycl_detail__::add_ir_attributes_global_variable(
306+
"sycl-device-global-size",
307+
__SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::name..., sizeof(T),
308+
__SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::value...)]]
309+
#endif
310+
device_global<
311+
T, detail::properties_t<Props...>,
312+
std::enable_if_t<!(detail::properties_t<Props...>::
313+
template has_property<device_constant_key>())>>
314+
: public detail::device_global_base<T, detail::properties_t<Props...>> {
241315

242-
device_global(const device_global &&) = delete;
243-
device_global &operator=(const device_global &) = delete;
244-
device_global &operator=(const device_global &&) = delete;
316+
using property_list_t = detail::properties_t<Props...>;
245317

246-
T &get() noexcept {
247-
__SYCL_HOST_NOT_SUPPORTED("get()")
248-
return *this->get_ptr();
249-
}
318+
public:
319+
using element_type = std::remove_extent_t<T>;
250320

251-
const T &get() const noexcept {
321+
#if !__cpp_consteval
322+
static_assert(std::is_trivially_default_constructible_v<T>,
323+
"Type T must be trivially default constructable (until C++20 "
324+
"consteval is supported and enabled.)");
325+
#endif // !__cpp_consteval
326+
static_assert(std::is_trivially_destructible_v<T>,
327+
"Type T must be trivially destructible.");
328+
329+
static_assert(is_property_list<property_list_t>::value,
330+
"Property list is invalid.");
331+
332+
DEVICE_GLOBAL_COMMON_METHODS()
333+
334+
T &get() noexcept {
252335
__SYCL_HOST_NOT_SUPPORTED("get()")
253336
return *this->get_ptr();
254337
}
@@ -258,58 +341,11 @@ class
258341
return get();
259342
}
260343

261-
operator const T &() const noexcept {
262-
__SYCL_HOST_NOT_SUPPORTED("Implicit conversion of device_global to T")
263-
return get();
264-
}
265-
266344
device_global &operator=(const T &newValue) noexcept {
267345
__SYCL_HOST_NOT_SUPPORTED("Assignment operator")
268346
*this->get_ptr() = newValue;
269347
return *this;
270348
}
271-
272-
template <class RelayT = T>
273-
std::remove_reference_t<
274-
decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> &
275-
operator[](std::ptrdiff_t idx) noexcept {
276-
__SYCL_HOST_NOT_SUPPORTED("Subscript operator")
277-
return (*this->get_ptr())[idx];
278-
}
279-
280-
template <class RelayT = T>
281-
const std::remove_reference_t<
282-
decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> &
283-
operator[](std::ptrdiff_t idx) const noexcept {
284-
__SYCL_HOST_NOT_SUPPORTED("Subscript operator")
285-
return (*this->get_ptr())[idx];
286-
}
287-
288-
template <class RelayT = T>
289-
std::enable_if_t<detail::HasArrowOperator<RelayT>::value ||
290-
std::is_pointer_v<RelayT>,
291-
RelayT> &
292-
operator->() noexcept {
293-
__SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global")
294-
return *this->get_ptr();
295-
}
296-
297-
template <class RelayT = T>
298-
std::enable_if_t<detail::HasArrowOperator<RelayT>::value ||
299-
std::is_pointer_v<RelayT>,
300-
const RelayT> &
301-
operator->() const noexcept {
302-
__SYCL_HOST_NOT_SUPPORTED("operator-> on a device_global")
303-
return *this->get_ptr();
304-
}
305-
306-
template <typename propertyT> static constexpr bool has_property() {
307-
return property_list_t::template has_property<propertyT>();
308-
}
309-
310-
template <typename propertyT> static constexpr auto get_property() {
311-
return property_list_t::template get_property<propertyT>();
312-
}
313349
};
314350

315351
} // namespace ext::oneapi::experimental

‎sycl/include/sycl/ext/oneapi/device_global/properties.hpp‎

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@
1010

1111
#include <sycl/ext/oneapi/properties/property.hpp> // for PropKind
1212
#include <sycl/ext/oneapi/properties/property_value.hpp> // for property_value
13+
#include <sycl/ext/oneapi/properties/properties.hpp> // for empty_properties_t
1314

1415
#include <cstdint> // for uint16_t
1516
#include <iosfwd> // for nullptr_t
@@ -19,13 +20,18 @@ namespace sycl {
1920
inline namespace _V1 {
2021
namespace ext::oneapi::experimental {
2122

22-
template <typename T, typename PropertyListT> class device_global;
23+
template <typename T, typename PropertyListT = empty_properties_t, typename = void> class device_global;
2324

2425
struct device_image_scope_key
2526
: detail::compile_time_property_key<detail::PropKind::DeviceImageScope> {
2627
using value_t = property_value<device_image_scope_key>;
2728
};
2829

30+
struct device_constant_key {
31+
: detail::compile_time_property_key<detail::PropKind::DeviceConstant> {
32+
using value_t = property_value<device_constant_key>;
33+
};
34+
2935
enum class host_access_enum : std::uint16_t { read, write, read_write, none };
3036

3137
struct host_access_key
@@ -54,6 +60,7 @@ struct implement_in_csr_key
5460
};
5561

5662
inline constexpr device_image_scope_key::value_t device_image_scope;
63+
inline constexpr device_constant_key::value_t device_constant;
5764

5865
template <host_access_enum Access>
5966
inline constexpr host_access_key::value_t<Access> host_access;
@@ -81,6 +88,9 @@ template <typename T, typename PropertyListT>
8188
struct is_property_key_of<device_image_scope_key,
8289
device_global<T, PropertyListT>> : std::true_type {};
8390
template <typename T, typename PropertyListT>
91+
struct is_property_key_of<device_constant_key,
92+
device_global<T, PropertyListT>> : std::true_type {};
93+
template <typename T, typename PropertyListT>
8494
struct is_property_key_of<host_access_key, device_global<T, PropertyListT>>
8595
: std::true_type {};
8696
template <typename T, typename PropertyListT>

‎sycl/include/sycl/ext/oneapi/properties/property.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -211,8 +211,9 @@ enum PropKind : uint32_t {
211211
Prefetch = 76,
212212
Deterministic = 77,
213213
InitializeToIdentity = 78,
214+
DeviceConstant= 79,
214215
// PropKindSize must always be the last value.
215-
PropKindSize = 79,
216+
PropKindSize = 80,
216217
};
217218

218219
struct property_key_base_tag {};

‎sycl/test-e2e/DeviceGlobal/device_constant.cpp‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
//
4+
// The OpenCL GPU backends do not currently support device_global backend
5+
// calls.
6+
//
7+
// UNSUPPORTED: opencl && gpu
8+
9+
#include "common.hpp"
10+
11+
device_global<int, decltype(properties{device_constant})> DeviceGlobalVar;
12+
13+
int main() {
14+
queue Q;
15+
16+
int HostVal = 42;
17+
Q.memcpy(DeviceGlobalVar, &HostVal);
18+
Q.wait();
19+
int OutVal = 0;
20+
21+
{
22+
buffer<int, 1> OutBuf(&OutVal, 1);
23+
Q.submit([&](handler &CGH) {
24+
auto OutAcc = OutBuf.get_access<access::mode::write>(CGH);
25+
CGH.single_task([=]() { OutAcc[0] = DeviceGlobalVar.get(); });
26+
});
27+
}
28+
assert(OutVal == 42 && "Read value does not match.");
29+
return 0;
30+
}

‎sycl/test/check_device_code/device_global_const.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
using namespace sycl;
99
using namespace sycl::ext::oneapi::experimental;
1010

11-
device_global<const int> DeviceGlobalVar;
11+
device_global<const int, decltype(properties{device_constant})> DeviceGlobalVar;
1212

1313
int main() {
1414
queue Q;

0 commit comments

Comments
 (0)