[SYCL] Ignore host_access attribute for non-device_image_scope device_global (#11194)

This commit makes a `device_global` without the `device_image_scope`
property ignore the `host_access` property in the kernel information, as
the runtime will need to access the `device_global` to set the pointer
inside and as such cannot respect the related promise. The headers will
still consider the `device_global` as having the property.

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

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

@@ -28,6 +28,14 @@
       Op " is not supported on host device.");
 #endif
 
+// Helper macro for conditional device_global property meta info filtering. This
+// lets us ignore certain properties under specified conditions, e.g. ignoring
+// host_access if device_image_scope isn't also present.
+#define __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)                             \
+  detail::ConditionalPropertyMetaInfo<                                         \
+      Props, detail::DeviceGlobalMetaInfoFilter<                               \
+                 Props, detail::properties_t<Props...>>::value>
+
 namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi::experimental {
@@ -120,8 +128,9 @@ class
 #ifdef __SYCL_DEVICE_ONLY__
     [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global,
       __sycl_detail__::add_ir_attributes_global_variable(
-          "sycl-device-global-size", detail::PropertyMetaInfo<Props>::name...,
-          sizeof(T), detail::PropertyMetaInfo<Props>::value...)]]
+          "sycl-device-global-size",
+          __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::name..., sizeof(T),
+          __SYCL_DEVICE_GLOBAL_PROP_META_INFO(Props)::value...)]]
 #endif
     device_global<T, detail::properties_t<Props...>>
     : public detail::device_global_base<T, detail::properties_t<Props...>> {
@@ -176,16 +185,16 @@ class
 
   template <class RelayT = T>
   std::remove_reference_t<
-      decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])>
-      &operator[](std::ptrdiff_t idx) noexcept {
+      decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> &
+  operator[](std::ptrdiff_t idx) noexcept {
     __SYCL_HOST_NOT_SUPPORTED("Subscript operator")
     return (*this->get_ptr())[idx];
   }
 
   template <class RelayT = T>
   const std::remove_reference_t<
-      decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])>
-      &operator[](std::ptrdiff_t idx) const noexcept {
+      decltype(std::declval<RelayT>()[std::declval<std::ptrdiff_t>()])> &
+  operator[](std::ptrdiff_t idx) const noexcept {
     __SYCL_HOST_NOT_SUPPORTED("Subscript operator")
     return (*this->get_ptr())[idx];
   }
@@ -222,3 +231,4 @@ class
 } // namespace sycl
 
 #undef __SYCL_HOST_NOT_SUPPORTED
+#undef __SYCL_DEVICE_GLOBAL_PROP_META_INFO

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

@@ -132,6 +132,19 @@ struct PropertyMetaInfo<implement_in_csr_key::value_t<Enable>> {
   static constexpr bool value = Enable;
 };
 
+// Filter allowing additional conditions for selecting when to include meta
+// information for properties for device_global.
+template <typename PropT, typename Properties>
+struct DeviceGlobalMetaInfoFilter : std::true_type {};
+
+// host_access cannot be honored for device_global variables without the
+// device_image_scope property, as the runtime needs to write the common USM
+// pointer during first launch.
+template <host_access_enum Access, typename Properties>
+struct DeviceGlobalMetaInfoFilter<host_access_key::value_t<Access>, Properties>
+    : std::bool_constant<
+          Properties::template has_property<device_image_scope_key>()> {};
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1

sycl/include/sycl/ext/oneapi/properties/property_utils.hpp

@@ -342,6 +342,24 @@ template <typename PropertiesT>
 struct NoConflictingProperties
     : NoConflictingPropertiesHelper<PropertiesT, PropertiesT> {};
 
+//******************************************************************************
+// Conditional property meta-info
+//******************************************************************************
+
+// Base class for property meta info that is ignored when propagating
+// information through the compiler.
+struct IgnoredPropertyMetaInfo {
+  static constexpr const char *name = "";
+  static constexpr std::nullptr_t value = nullptr;
+};
+
+// Trait for picking either property meta information for PropT if Condition is
+// true or ignored information if Condition is false.
+template <typename PropT, bool Condition>
+struct ConditionalPropertyMetaInfo
+    : std::conditional_t<Condition, PropertyMetaInfo<PropT>,
+                         IgnoredPropertyMetaInfo> {};
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1

sycl/test/extensions/device_global/device_global_ignore_host_access.cpp

@@ -0,0 +1,27 @@
+// RUN: %clangxx -S -emit-llvm -fsycl-device-only %s -o - | FileCheck %s
+
+// Checks that the host_access property doesn't get represented when there is no
+// device_image_scope property.
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::oneapi::experimental;
+
+static device_global<int,
+                     decltype(properties(device_image_scope, host_access_read))>
+    DeviceGlobalDeviceImageScoped;
+static device_global<int, decltype(properties(host_access_read))>
+    DeviceGlobalFullScoped;
+
+int main() {
+  sycl::queue Q;
+  Q.single_task([]() {
+    DeviceGlobalFullScoped = 42;
+    DeviceGlobalDeviceImageScoped = 42;
+  });
+}
+
+// CHECK-DAG: @_ZL29DeviceGlobalDeviceImageScoped = {{.*}} #[[DISAttrs:[0-9]+]]
+// CHECK-DAG: @_ZL22DeviceGlobalFullScoped = {{.*}} #[[FSAttrs:[0-9]+]]
+// CHECK-DAG: attributes #[[DISAttrs:[0-9]+]] = { {{.*}}"sycl-host-access"
+// CHECK-NOT: attributes #[[FSAttrs:[0-9]+]] = { {{.*}}"sycl-host-access"