[SYCL] Introduce new properties for virtual functions (#14014)

See language specification for those properties in #10540

Author: AlexeySachkov

sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp

@@ -0,0 +1,74 @@
+#pragma once
+
+#include <sycl/ext/oneapi/properties/property.hpp>
+#include <sycl/ext/oneapi/properties/property_value.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+struct indirectly_callable_key {
+  template <typename Set>
+  using value_t =
+      sycl::ext::oneapi::experimental::property_value<indirectly_callable_key,
+                                                      Set>;
+};
+
+template <typename Set = void>
+inline constexpr indirectly_callable_key::value_t<Set> indirectly_callable;
+
+struct calls_indirectly_key {
+  template <typename First = void, typename... SetIds>
+  using value_t =
+      sycl::ext::oneapi::experimental::property_value<calls_indirectly_key,
+                                                      First, SetIds...>;
+};
+
+template <typename First = void, typename... Rest>
+inline constexpr calls_indirectly_key::value_t<First, Rest...> calls_indirectly;
+
+template <> struct is_property_key<indirectly_callable_key> : std::true_type {};
+template <> struct is_property_key<calls_indirectly_key> : std::true_type {};
+
+namespace detail {
+
+template <>
+struct IsCompileTimeProperty<indirectly_callable_key> : std::true_type {};
+template <>
+struct IsCompileTimeProperty<calls_indirectly_key> : std::true_type {};
+
+template <> struct PropertyToKind<indirectly_callable_key> {
+  static constexpr PropKind Kind = PropKind::IndirectlyCallable;
+};
+
+template <> struct PropertyToKind<calls_indirectly_key> {
+  static constexpr PropKind Kind = PropKind::CallsIndirectly;
+};
+
+template <typename Set>
+struct PropertyMetaInfo<indirectly_callable_key::value_t<Set>> {
+  static constexpr const char *name = "indirectly-callable";
+  static constexpr const char *value =
+#ifdef __SYCL_DEVICE_ONLY__
+      __builtin_sycl_unique_stable_name(Set);
+#else
+      "";
+#endif
+};
+
+template <typename First, typename... Rest>
+struct PropertyMetaInfo<calls_indirectly_key::value_t<First, Rest...>> {
+  static constexpr const char *name = "calls-indirectly";
+  static constexpr const char *value =
+#ifdef __SYCL_DEVICE_ONLY__
+      // FIXME: we should handle Rest... here as well
+      __builtin_sycl_unique_stable_name(First);
+#else
+      "";
+#endif
+};
+
+} // namespace detail
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

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

@@ -205,8 +205,10 @@ enum PropKind : uint32_t {
   WorkItemProgress = 64,
   NDRangeKernel = 65,
   SingleTaskKernel = 66,
+  IndirectlyCallable = 67,
+  CallsIndirectly = 68,
   // PropKindSize must always be the last value.
-  PropKindSize = 67,
+  PropKindSize = 69,
 };
 
 struct property_key_base_tag {};

sycl/include/sycl/handler.hpp

@@ -34,6 +34,7 @@
 #include <sycl/ext/oneapi/device_global/properties.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
+#include <sycl/ext/oneapi/experimental/virtual_functions.hpp>
 #include <sycl/ext/oneapi/kernel_properties/properties.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/group.hpp>
@@ -962,6 +963,10 @@ class __SYCL_EXPORT handler {
                  sycl::ext::intel::experimental::fp_control_key>() &&
              KI::isESIMD()),
         "Floating point control property is supported for ESIMD kernels only.");
+    static_assert(
+        !PropertiesT::template has_property<
+            sycl::ext::oneapi::experimental::indirectly_callable_key>(),
+        "indirectly_callable property cannot be applied to SYCL kernels");
     if constexpr (PropertiesT::template has_property<
                       sycl::ext::intel::experimental::cache_config_key>()) {
       auto Config = Props.template get_property<

sycl/test/include_deps/sycl_detail_core.hpp.cpp

@@ -178,5 +178,6 @@
 // CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
 // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
 // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
 // CHECK-EMPTY:

sycl/test/virtual-functions/calls-indirectly-ir.cpp

@@ -0,0 +1,57 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -emit-llvm -S %s -o %t.ll
+// RUN: FileCheck %s < %t.ll
+//
+// This test is intended to check integration between SYCL headers and SYCL FE,
+// i.e. to make sure that setting properties related to virtual functions will
+// result in the right LLVM IR.
+//
+// This test is specifically focused on the calls_indirectly property.
+//
+// CHECK: define {{.*}}KEmpty{{.*}} #[[#ATTR_SET_DEFAULT:]]
+// CHECK: define {{.*}}KInt{{.*}} #[[#ATTR_SET_INT:]]
+// CHECK: define {{.*}}KVoid{{.*}} #[[#ATTR_SET_DEFAULT]]
+// CHECK: define {{.*}}KUserDefined{{.*}} #[[#ATTR_SET_USER_DEFINED:]]
+// TODO: update the check below
+// As of now calls_indirectly_property takes into account only the first
+// template argument ignoring the rest. This will be fixed in a follow-up
+// patches and the test should be updated to reflect that, because current
+// behavior is not correct.
+// CHECK: define {{.*}}KMultiple{{.*}} #[[#ATTR_SET_INT]]
+//
+// CHECK-DAG: attributes #[[#ATTR_SET_DEFAULT]] {{.*}} "calls-indirectly"="_ZTSv"
+// CHECK-DAG: attributes #[[#ATTR_SET_INT]] {{.*}} "calls-indirectly"="_ZTSi"
+// CHECK-DAG: attributes #[[#ATTR_SET_USER_DEFINED]] {{.*}} "calls-indirectly"="_ZTS12user_defined"
+
+#include <sycl/sycl.hpp>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+struct user_defined {
+  int a;
+  float b;
+};
+
+class KEmpty;
+class KInt;
+class KVoid;
+class KUserDefined;
+class KMultiple;
+
+int main() {
+  sycl::queue q;
+
+  oneapi::properties props_empty{oneapi::calls_indirectly<>};
+  oneapi::properties props_int{oneapi::calls_indirectly<int>};
+  oneapi::properties props_void{oneapi::calls_indirectly<void>};
+  oneapi::properties props_user_defined{oneapi::calls_indirectly<user_defined>};
+  oneapi::properties props_multiple{
+      oneapi::calls_indirectly<int, user_defined>};
+
+  q.single_task<KEmpty>(props_empty, [=]() {});
+  q.single_task<KInt>(props_int, [=]() {});
+  q.single_task<KVoid>(props_void, [=]() {});
+  q.single_task<KUserDefined>(props_user_defined, [=]() {});
+  q.single_task<KMultiple>(props_multiple, [=]() {});
+
+  return 0;
+}

sycl/test/virtual-functions/indirectly-callable-ir.cpp

@@ -0,0 +1,66 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -emit-llvm -S -Xclang -fsycl-allow-virtual-functions %s -o %t.ll
+// RUN: FileCheck %s < %t.ll
+//
+// This test is intended to check integration between SYCL headers and SYCL FE,
+// i.e. to make sure that setting properties related to virtual functions will
+// result in the right LLVM IR.
+//
+// This test is specifically focused on the indirectly_callable property.
+//
+// CHECK: define {{.*}} @_ZN4Base3fooEv{{.*}} #[[#ATTR_SET_DEFAULT:]]
+// CHECK: define {{.*}} @_ZN7Derived3fooEv{{.*}} #[[#ATTR_SET_DEFAULT]]
+// CHECK: define {{.*}} @_ZN7Derived3barEv{{.*}} #[[#ATTR_SET_DEFAULT]]
+// CHECK: define {{.*}} @_ZN10SubDerived3barEv{{.*}} #[[#ATTR_SET_INT:]]
+// CHECK: define {{.*}} @_ZN13SubSubDerived3foo{{.*}} #[[#ATTR_SET_DEFAULT]]
+// CHECK: define {{.*}} @_ZN13SubSubDerived3barEv{{.*}} #[[#ATTR_SET_BASE:]]
+//
+// CHECK-DAG: attributes #[[#ATTR_SET_DEFAULT]] {{.*}} "indirectly-callable"="_ZTSv"
+// CHECK-DAG: attributes #[[#ATTR_SET_INT]] {{.*}} "indirectly-callable"="_ZTSi"
+// CHECK-DAG: attributes #[[#ATTR_SET_BASE]] {{.*}} "indirectly-callable"="_ZTS4Base"
+
+#include <sycl/sycl.hpp>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual int foo();
+};
+
+int Base::foo() { return 42; }
+
+class Derived : public Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<void>)
+  virtual int bar();
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  int foo() override;
+};
+
+int Derived::foo() { return 43; }
+
+int Derived::bar() { return 0; }
+
+class SubDerived : public Derived {
+public:
+  int foo() override { return 44; }
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<int>)
+  int bar() override;
+};
+
+int SubDerived::bar() { return 1; }
+
+class SubSubDerived : public SubDerived {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  int foo() override;
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<Base>)
+  int bar() override;
+};
+
+int SubSubDerived::foo() { return 45; }
+
+int SubSubDerived::bar() { return 2; }

sycl/test/virtual-functions/properties-negative.cpp

@@ -0,0 +1,30 @@
+// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
+
+#include <sycl/sycl.hpp>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+struct user_defined {
+  int a;
+  float b;
+};
+
+int main() {
+  sycl::queue q;
+
+  oneapi::properties props_empty{oneapi::indirectly_callable<>};
+  oneapi::properties props_void{oneapi::indirectly_callable<void>};
+  oneapi::properties props_int{oneapi::indirectly_callable<int>};
+  oneapi::properties props_user{oneapi::indirectly_callable<user_defined>};
+
+  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  q.single_task(props_empty, [=]() {});
+  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  q.single_task(props_void, [=]() {});
+  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  q.single_task(props_int, [=]() {});
+  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  q.single_task(props_user, [=]() {});
+
+  return 0;
+}

sycl/test/virtual-functions/properties-positive.cpp

@@ -0,0 +1,64 @@
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
+//
+// This test is intended to check that we can successfully compile code that
+// uses new properties from the virtual functions extension.
+//
+// expected-no-diagnostics
+
+#include <sycl/sycl.hpp>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+class Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  virtual void foo() {}
+};
+
+class Derived : public Base {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void foo() override {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<void>)
+  virtual void bar() {}
+};
+
+class SubDerived : public Derived {
+public:
+  void foo() override {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<int>)
+  void bar() override {}
+};
+
+class SubSubDerived : public SubDerived {
+public:
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<>)
+  void foo() override {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable<Base>)
+  void bar() override {}
+};
+
+int main() {
+  sycl::queue q;
+
+  static_assert(
+      oneapi::is_property_key<oneapi::indirectly_callable_key>::value);
+  static_assert(oneapi::is_property_key<oneapi::calls_indirectly_key>::value);
+
+  oneapi::properties props_empty{oneapi::calls_indirectly<>};
+  oneapi::properties props_void{oneapi::calls_indirectly<void>};
+  oneapi::properties props_int{oneapi::calls_indirectly<int>};
+  oneapi::properties props_base{oneapi::calls_indirectly<Base>};
+  oneapi::properties props_multiple{oneapi::calls_indirectly<int, Base>};
+
+  q.single_task(props_empty, [=]() {});
+  q.single_task(props_void, [=]() {});
+  q.single_task(props_int, [=]() {});
+  q.single_task(props_base, [=]() {});
+  q.single_task(props_multiple, [=]() {});
+
+  return 0;
+}