[SYCL] Update calls_indirectly property (#15523)

This allows calls_indirectly property to handle all passed template
arguments, but not the only first one.

Also applying some minor changes to the spec.

Spec & Design: #10540

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

@@ -19,16 +19,16 @@ template <typename Set>
 inline constexpr indirectly_callable_key::value_t<Set> indirectly_callable_in;
 
 struct calls_indirectly_key {
-  template <typename First = void, typename... SetIds>
+  template <typename... SetIds>
   using value_t =
       sycl::ext::oneapi::experimental::property_value<calls_indirectly_key,
-                                                      First, SetIds...>;
+                                                      SetIds...>;
 };
 
 inline constexpr calls_indirectly_key::value_t<void> assume_indirect_calls;
 
-template <typename First, typename... Rest>
-inline constexpr calls_indirectly_key::value_t<First, Rest...>
+template <typename... SetIds>
+inline constexpr calls_indirectly_key::value_t<SetIds...>
     assume_indirect_calls_to;
 
 template <> struct is_property_key<indirectly_callable_key> : std::true_type {};
@@ -60,16 +60,57 @@ struct PropertyMetaInfo<indirectly_callable_key::value_t<Set>> {
 #endif
 };
 
-template <typename First, typename... Rest>
-struct PropertyMetaInfo<calls_indirectly_key::value_t<First, Rest...>> {
-  static_assert(
-      sizeof...(Rest) == 0,
-      "assume_indirect_calls_to property only supports a single set for now");
+#ifdef __SYCL_DEVICE_ONLY__
+// Helper to concatenate several lists of characters into a single string.
+// Lists are separated from each other with comma within the resulting string.
+template <typename List, typename... Rest> struct ConcatenateCharsToStr;
+
+// Specialization for a single list
+template <char... Chars> struct ConcatenateCharsToStr<CharList<Chars...>> {
+  static constexpr char value[] = {Chars..., '\0'};
+};
+
+// Specialization for two lists
+template <char... Chars, char... CharsToAppend>
+struct ConcatenateCharsToStr<CharList<Chars...>, CharList<CharsToAppend...>>
+    : ConcatenateCharsToStr<CharList<Chars..., ',', CharsToAppend...>> {};
+
+// Specialization for the case when there are more than two lists
+template <char... Chars, char... CharsToAppend, typename... Rest>
+struct ConcatenateCharsToStr<CharList<Chars...>, CharList<CharsToAppend...>,
+                             Rest...>
+    : ConcatenateCharsToStr<CharList<Chars..., ',', CharsToAppend...>,
+                            Rest...> {};
+
+// Helper to convert type T to a list of characters representing the type (its
+// mangled name).
+template <typename T, size_t... Indices> struct StableNameToCharsHelper {
+  using chars = CharList<__builtin_sycl_unique_stable_name(T)[Indices]...>;
+};
+
+// Wrapper helper for the struct above
+template <typename T, typename Sequence> struct StableNameToChars;
+
+// Specialization of that wrapper helper which accepts sequence of integers
+template <typename T, size_t... Indices>
+struct StableNameToChars<T, std::integer_sequence<size_t, Indices...>>
+    : StableNameToCharsHelper<T, Indices...> {};
+
+// Creates a comma-separated string with unique stable names for each type in
+// Ts.
+template <typename... Ts>
+struct UniqueStableNameListStr
+    : ConcatenateCharsToStr<typename StableNameToChars<
+          Ts, std::make_index_sequence<__builtin_strlen(
+                  __builtin_sycl_unique_stable_name(Ts))>>::chars...> {};
+#endif // __SYCL_DEVICE_ONLY__
+
+template <typename... SetIds>
+struct PropertyMetaInfo<calls_indirectly_key::value_t<SetIds...>> {
   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);
+      UniqueStableNameListStr<SetIds...>::value;
 #else
       "";
 #endif

sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp

@@ -51,9 +51,18 @@ class IncrementBy8 : public BaseIncrement {
   void increment(int *Data) override { *Data += 8 + Mod; }
 };
 
+struct SetIncBy16;
+class IncrementBy16 : public BaseIncrement {
+public:
+  IncrementBy16(int Mod, int /* unused */) : BaseIncrement(Mod) {}
+
+  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(oneapi::indirectly_callable_in<SetIncBy16>)
+  void increment(int *Data) override { *Data += 16 + Mod; }
+};
+
 int main() try {
-  using storage_t =
-      obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4, IncrementBy8>;
+  using storage_t = obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4,
+                                  IncrementBy8, IncrementBy16>;
 
   storage_t HostStorage;
   sycl::buffer<storage_t> DeviceStorage(sycl::range{1});
@@ -66,8 +75,9 @@ int main() try {
   sycl::queue q(asyncHandler);
 
   // TODO: cover uses case when objects are passed through USM
-  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
-  for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
+  constexpr oneapi::properties props{
+      oneapi::assume_indirect_calls_to<void, SetIncBy16>};
+  for (unsigned TestCase = 0; TestCase < 5; ++TestCase) {
     int HostData = 42;
     int Data = HostData;
     sycl::buffer<int> DataStorage(&Data, sycl::range{1});

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

@@ -11,16 +11,12 @@
 // 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-disabled: define {{.*}}KMultiple{{.*}} #[[#ATTR_SET_INT]]
+// CHECK: define {{.*}}KMultiple{{.*}} #[[#ATTR_SET_MULTIPLE:]]
 //
 // 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"
+// CHECK-DAG: attributes #[[#ATTR_SET_MULTIPLE]] {{.*}} "calls-indirectly"="_ZTSi,_ZTS12user_defined"
 
 #include <sycl/sycl.hpp>
 
@@ -45,17 +41,14 @@ int main() {
   oneapi::properties props_void{oneapi::assume_indirect_calls_to<void>};
   oneapi::properties props_user_defined{
       oneapi::assume_indirect_calls_to<user_defined>};
-  // assume_indirect_calls_to is currently limited to a single set, so this test
-  // is disabled.
-  // FIXME: re-enable once the restriction is lifted.
-  // oneapi::properties props_multiple{
-  //    oneapi::assume_indirect_calls_to<int, user_defined>};
+  oneapi::properties props_multiple{
+     oneapi::assume_indirect_calls_to<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, [=]() {});
+  q.single_task<KMultiple>(props_multiple, [=]() {});
 
   return 0;
 }

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

@@ -52,17 +52,14 @@ int main() {
   oneapi::properties props_void{oneapi::assume_indirect_calls_to<void>};
   oneapi::properties props_int{oneapi::assume_indirect_calls_to<int>};
   oneapi::properties props_base{oneapi::assume_indirect_calls_to<Base>};
-  // assume_indirect_calls_to is currently limited to a single set, so this test
-  // is disabled.
-  // FIXME: re-enable once the restriction is lifted.
-  // oneapi::properties props_multiple{
-  //    oneapi::assume_indirect_calls_to<int, Base>};
+  oneapi::properties props_multiple{
+     oneapi::assume_indirect_calls_to<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, [=]() {});
+  q.single_task(props_multiple, [=]() {});
 
   return 0;
 }