diff --git a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp index acdb8c1ea25ae..a3100dc1a2657 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/virtual_functions.hpp @@ -19,16 +19,16 @@ template inline constexpr indirectly_callable_key::value_t indirectly_callable_in; struct calls_indirectly_key { - template + template using value_t = sycl::ext::oneapi::experimental::property_value; + SetIds...>; }; inline constexpr calls_indirectly_key::value_t assume_indirect_calls; -template -inline constexpr calls_indirectly_key::value_t +template +inline constexpr calls_indirectly_key::value_t assume_indirect_calls_to; template <> struct is_property_key : std::true_type {}; @@ -60,16 +60,57 @@ struct PropertyMetaInfo> { #endif }; -template -struct PropertyMetaInfo> { - 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 struct ConcatenateCharsToStr; + +// Specialization for a single list +template struct ConcatenateCharsToStr> { + static constexpr char value[] = {Chars..., '\0'}; +}; + +// Specialization for two lists +template +struct ConcatenateCharsToStr, CharList> + : ConcatenateCharsToStr> {}; + +// Specialization for the case when there are more than two lists +template +struct ConcatenateCharsToStr, CharList, + Rest...> + : ConcatenateCharsToStr, + Rest...> {}; + +// Helper to convert type T to a list of characters representing the type (its +// mangled name). +template struct StableNameToCharsHelper { + using chars = CharList<__builtin_sycl_unique_stable_name(T)[Indices]...>; +}; + +// Wrapper helper for the struct above +template struct StableNameToChars; + +// Specialization of that wrapper helper which accepts sequence of integers +template +struct StableNameToChars> + : StableNameToCharsHelper {}; + +// Creates a comma-separated string with unique stable names for each type in +// Ts. +template +struct UniqueStableNameListStr + : ConcatenateCharsToStr>::chars...> {}; +#endif // __SYCL_DEVICE_ONLY__ + +template +struct PropertyMetaInfo> { 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::value; #else ""; #endif diff --git a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp b/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp index a2cea9a6a0120..5f8c7c9323465 100644 --- a/sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp +++ b/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) + void increment(int *Data) override { *Data += 16 + Mod; } +}; + int main() try { - using storage_t = - obj_storage_t; + using storage_t = obj_storage_t; storage_t HostStorage; sycl::buffer 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}; + for (unsigned TestCase = 0; TestCase < 5; ++TestCase) { int HostData = 42; int Data = HostData; sycl::buffer DataStorage(&Data, sycl::range{1}); diff --git a/sycl/test/virtual-functions/calls-indirectly-ir.cpp b/sycl/test/virtual-functions/calls-indirectly-ir.cpp index 89b238d895728..2857cc21d5c03 100644 --- a/sycl/test/virtual-functions/calls-indirectly-ir.cpp +++ b/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 @@ -45,17 +41,14 @@ int main() { oneapi::properties props_void{oneapi::assume_indirect_calls_to}; oneapi::properties props_user_defined{ oneapi::assume_indirect_calls_to}; - // 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}; + oneapi::properties props_multiple{ + oneapi::assume_indirect_calls_to}; q.single_task(props_empty, [=]() {}); q.single_task(props_int, [=]() {}); q.single_task(props_void, [=]() {}); q.single_task(props_user_defined, [=]() {}); - // q.single_task(props_multiple, [=]() {}); + q.single_task(props_multiple, [=]() {}); return 0; } diff --git a/sycl/test/virtual-functions/properties-positive.cpp b/sycl/test/virtual-functions/properties-positive.cpp index e4ebc06601804..970e348457388 100644 --- a/sycl/test/virtual-functions/properties-positive.cpp +++ b/sycl/test/virtual-functions/properties-positive.cpp @@ -52,17 +52,14 @@ int main() { oneapi::properties props_void{oneapi::assume_indirect_calls_to}; oneapi::properties props_int{oneapi::assume_indirect_calls_to}; oneapi::properties props_base{oneapi::assume_indirect_calls_to}; - // 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}; + oneapi::properties props_multiple{ + oneapi::assume_indirect_calls_to}; 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; }