From 071702630125e63e822cdc7d01c56a0babebdf96 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 3 Mar 2025 09:45:43 +0000 Subject: [PATCH 1/2] [SYCL][RTC] Rename property to `registered_names` Signed-off-by: Julian Oppermann --- .../sycl_ext_oneapi_kernel_compiler.asciidoc | 82 +++++++++---------- .../sycl/ext/oneapi/properties/property.hpp | 2 +- sycl/include/sycl/kernel_bundle.hpp | 27 +++--- .../KernelCompiler/kernel_compiler_sycl.cpp | 2 +- .../kernel_compiler_sycl_jit.cpp | 8 +- .../abi/sycl_classes_abi_neutral_test.cpp | 4 +- 6 files changed, 61 insertions(+), 64 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index d796e49b628a7..c0291e6a20ede 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -485,53 +485,52 @@ a! ---- namespace sycl::ext::oneapi::experimental { -struct registered_kernel_names { - registered_kernel_names(); (1) - registered_kernel_names(const std::string &name); (2) - registered_kernel_names(const std::vector &names); (3) - void add(const std::string &name); (4) +struct registered_names { + registered_names(); (1) + registered_names(const std::string &name); (2) + registered_names(const std::vector &names); (3) + void add(const std::string &name); (4) }; -using registered_kernel_names_key = registered_kernel_names; +using registered_names_key = registered_names; template<> -struct is_property_key : std::true_type {}; +struct is_property_key : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ---- !==== -This property is useful when the source language represents kernel names -differently in the source code and the generated code. +This property is useful when the source language represents names differently in +the source code and the generated code. For example, {cpp} function names in the generated code are "mangled" in an implementation-defined way. The precise meaning of this property is defined by each source language, but in -general it allows the application to supply a list of kernel names as they -appear in the source code. +general it allows the application to supply a list of names as they appear in +the source code. The application can then get the corresponding raw (i.e. mangled) names after the code is compiled. See the section below "Obtaining a kernel when the language is ``sycl``" for a description of how this property is used with the `source_language::sycl` language. -_Effects (1):_ Creates a new `registered_kernel_names` property with no -registered kernel names. +_Effects (1):_ Creates a new `registered_names` property with no registered +names. -_Effects (2):_ Creates a new `registered_kernel_names` property with a single -registered kernel name. +_Effects (2):_ Creates a new `registered_names` property with a single +registered name. -_Effects (3):_ Creates a new `registered_kernel_names` property from a vector -of kernel names. +_Effects (3):_ Creates a new `registered_names` property from a vector of names. -_Effects (4):_ Adds `name` to the property's list of registered kernel names. +_Effects (4):_ Adds `name` to the property's list of registered names. -_Preconditions (2-4):_ Each source language defines its own requirements for -the registered kernel names. +_Preconditions (2-4):_ Each source language defines its own requirements for the +registered names. For the language `source_language::sycl`, each name must be a {cpp} expression for a pointer to a kernel function as defined below under "Obtaining a kernel when the language is ``sycl``". -[_Note:_ It is not an error to have duplicate names in a -`registered_kernel_names` property, but the duplicates have no effect. +[_Note:_ It is not an error to have duplicate names in a `registered_names` +property, but the duplicates have no effect. _{endnote}_] |==== @@ -643,8 +642,8 @@ _Constraints:_ This function is not available when `State` is _Returns:_ If the kernel bundle was created from a bundle of state `bundle_state::ext_oneapi_source` and `name` was registered via -`registered_kernel_names`, returns the compiler-generated (e.g. mangled) name -for this kernel function. +`registered_names`, returns the compiler-generated (e.g. mangled) name for this +kernel function. If the kernel bundle was created from a bundle of state `bundle_state::ext_oneapi_source` and `name` is the same as a compiler-generated name for a kernel defined in that bundle, that same @@ -661,7 +660,7 @@ _Throws:_ When the kernel is defined in the language `source_language::sycl`, the host code may query for the kernel or obtain the `kernel` object using either the kernel's name as it is generated by the compiler (i.e. the {cpp} mangled name) -or by using the `registered_kernel_names` property. +or by using the `registered_names` property. ==== Using the compiler-generated name @@ -691,11 +690,10 @@ sycl::kernel_bundle kb = /*...*/; sycl::kernel k = kb.ext_oneapi_get_kernel("foo"); ---- -==== Using the `registered_kernel_names` property +==== Using the `registered_names` property When the kernel is not declared as `extern "C"`, the compiler generates a -mangled name, so it is more convenient to use the `registered_kernel_names` -property. +mangled name, so it is more convenient to use the `registered_names` property. Each string in the property must be the {cpp} expression for a pointer to a kernel function. These expression strings are conceptually compiled at the bottom of source @@ -724,7 +722,7 @@ The host code can compile this and get the kernel's `kernel` object like so: sycl::kernel_bundle kb_src = /*...*/; sycl::kernel_bundle kb = syclexp::build(kb_src, - syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); + syclexp::properties{syclexp::registered_names{"mykernels::bar"}}); sycl::kernel k = kb.ext_oneapi_get_kernel("mykernels::bar"); ---- @@ -744,21 +742,21 @@ the kernel by calling `ext_oneapi_get_raw_kernel_name` like this: sycl::kernel_bundle kb_src = /*...*/; sycl::kernel_bundle kb = syclexp::build(kb_src, - syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); + syclexp::properties{syclexp::registered_names{"mykernels::bar"}}); std::string mangled_name = kb.ext_oneapi_get_raw_kernel_name("mykernels::bar"); ---- Again, the string passed to `ext_oneapi_get_raw_kernel_name` must have exactly -the same content as the string that was used to construct the -`registered_kernel_names` property. +the same content as the string that was used to construct the `registered_names` +property. The application may also pass this compiler-generated (i.e. mangled) name to `ext_oneapi_get_kernel` in order to get the `kernel` object. ==== Instantiating templated kernel functions -The `registered_kernel_names` property can also be used to instantiate a -kernel that is defined as a function template. +The `registered_names` property can also be used to instantiate a kernel that is +defined as a function template. For example, consider source code that defines a kernel function template like this: @@ -774,8 +772,8 @@ std::string source = R"""( )"""; ---- -The application can use the `registered_kernel_names` property to instantiate -the template for specific template arguments. +The application can use the `registered_names` property to instantiate the +template for specific template arguments. For example, this host code instantiates the template twice and gets a `kernel` object for each instantiation: @@ -784,7 +782,7 @@ object for each instantiation: sycl::kernel_bundle kb_src = /*...*/; sycl::kernel_bundle kb = syclexp::build(kb_src, - syclexp::properties{syclexp::registered_kernel_names{{"bartmpl", "bartmpl"}}); + syclexp::properties{syclexp::registered_names{{"bartmpl", "bartmpl"}}); sycl::kernel k_float = kb.ext_oneapi_get_kernel("bartmpl"); sycl::kernel k_int = kb.ext_oneapi_get_kernel("bartmpl"); @@ -830,7 +828,7 @@ int main() { syclexp::source_language::sycl, source); - // Compile the kernel. There is no need to use the "registered_kernel_names" + // Compile the kernel. There is no need to use the "registered_names" // property because the kernel is declared extern "C". sycl::kernel_bundle kb_exe = syclexp::build(kb_src); @@ -852,7 +850,7 @@ int main() { === Disambiguating overloaded kernel functions -This example demonstrates how to use the `registered_kernel_names` property to +This example demonstrates how to use the `registered_names` property to disambiguate a kernel function that has several overloads. [source,c++] @@ -896,11 +894,11 @@ int main() { // use a C++ cast to disambiguate between them. Here, we are selecting the // "int" overload. std::string iota_name{"(void(*)(int, int*))iota"}; - sycl::kernel_bundle kb_exe = syclexp::build(kb_src, - syclexp::properties{syclexp::registered_kernel_names{iota_name}}); + sycl::kernel_bundle kb_exe = + syclexp::build(kb_src, syclexp::properties{syclexp::registered_names{iota_name}}); // Get the kernel by passing the same string we used to construct the - // "registered_kernel_names" property. + // "registered_names" property. sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name); int *ptr = sycl::malloc_shared(NUM, q); diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 9fd7752889f4f..260c6ea82e7b0 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -210,7 +210,7 @@ enum PropKind : uint32_t { InputDataPlacement = 65, OutputDataPlacement = 66, IncludeFiles = 67, - RegisteredKernelNames = 68, + RegisteredNames = 68, ClusterLaunch = 69, FPGACluster = 70, Balanced = 71, diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index d77ef4319ed52..2fd4615020ebd 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -1001,22 +1001,21 @@ struct is_property_key_of : std::true_type {}; ///////////////////////// -// PropertyT syclex::registered_kernel_names +// PropertyT syclex::registered_names ///////////////////////// -struct registered_kernel_names - : detail::run_time_property_key { - std::vector kernel_names; - registered_kernel_names() {} - registered_kernel_names(const std::string &knArg) : kernel_names{knArg} {} - registered_kernel_names(const std::vector &knsArg) - : kernel_names(knsArg) {} - void add(const std::string &name) { kernel_names.push_back(name); } +struct registered_names + : detail::run_time_property_key { + std::vector names; + registered_names() {} + registered_names(const std::string &name) : names{name} {} + registered_names(const std::vector &names) : names{names} {} + void add(const std::string &name) { names.push_back(name); } }; -using registered_kernel_names_key = registered_kernel_names; +using registered_names_key = registered_names; template <> -struct is_property_key_of : std::true_type { }; @@ -1161,9 +1160,9 @@ build(kernel_bundle &SourceKB, if constexpr (props.template has_property()) { LogPtr = props.template get_property().log; } - if constexpr (props.template has_property()) { + if constexpr (props.template has_property()) { RegisteredKernelNamesVec = - props.template get_property().kernel_names; + props.template get_property().names; } return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr, RegisteredKernelNamesVec); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index eb4e95bbddfc1..848ffae1581ba 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -166,7 +166,7 @@ void test_build_and_run() { exe_kb kbExe2 = syclex::build( kbSrc, devs, syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}, - syclex::registered_kernel_names{"ff_templated"}}); + syclex::registered_names{"ff_templated"}}); assert(log.find("warning: 'this_nd_item<1>' is deprecated") != std::string::npos); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 67a80fdb55d0e..87037a63f3d39 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -236,7 +236,7 @@ int test_build_and_run() { exe_kb kbExe2 = syclex::build( kbSrc, devs, syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}, - syclex::registered_kernel_names{"ff_templated"}}); + syclex::registered_names{"ff_templated"}}); // extern "C" was used, so the name "ff_cp" is implicitly known. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); @@ -355,7 +355,7 @@ int test_device_code_split() { auto build = [&](const std::string &mode) -> size_t { exe_kb kbExe = syclex::build( kbSrc, syclex::properties{ - syclex::registered_kernel_names{names}, + syclex::registered_names{names}, syclex::build_options{"-fsycl-device-code-split=" + mode}}); return std::distance(kbExe.begin(), kbExe.end()); }; @@ -372,8 +372,8 @@ int test_device_code_split() { // Test implicit device code split names = {"vec_add", "vec_add"}; - exe_kb kbDiffWorkGroupSizes = syclex::build( - kbSrc, syclex::properties{syclex::registered_kernel_names{names}}); + exe_kb kbDiffWorkGroupSizes = + syclex::build(kbSrc, syclex::properties{syclex::registered_names{names}}); assert(std::distance(kbDiffWorkGroupSizes.begin(), kbDiffWorkGroupSizes.end()) == 2); diff --git a/sycl/test/abi/sycl_classes_abi_neutral_test.cpp b/sycl/test/abi/sycl_classes_abi_neutral_test.cpp index fc494d500ce90..9ecc45593f6de 100644 --- a/sycl/test/abi/sycl_classes_abi_neutral_test.cpp +++ b/sycl/test/abi/sycl_classes_abi_neutral_test.cpp @@ -33,9 +33,9 @@ // CHECK-NEXT: 0 | class {{(std::__new_allocator|__gnu_cxx::new_allocator)}}, class std::basic_string > > (base) (empty) // CHECK-NEXT: 0 | {{(struct std::_Vector_base, class std::basic_string >, class std::allocator, class std::basic_string > > >::_Vector_impl_data \(base\)|pointer _M_start)}} -// CHECK: 0 | struct sycl::ext::oneapi::experimental::registered_kernel_names +// CHECK: 0 | struct sycl::ext::oneapi::experimental::registered_names // CHECK-NEXT: 0 | struct sycl::ext::oneapi::experimental::detail::run_time_property_key -// CHECK: 0 | class std::vector > kernel_names +// CHECK: 0 | class std::vector > names // CHECK-NEXT: 0 | struct std::_Vector_base, class std::allocator > > (base) // CHECK-NEXT: 0 | struct std::_Vector_base, class std::allocator > >::_Vector_impl _M_impl // CHECK-NEXT: 0 | class std::allocator > (base) (empty) From 43ad8c547456d7d0060598b766ee1cb9b14a0e32 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Mon, 3 Mar 2025 21:11:03 +0000 Subject: [PATCH 2/2] Second example for mangled names Signed-off-by: Julian Oppermann --- .../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index c0291e6a20ede..86689b94cae69 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -502,8 +502,8 @@ struct is_property_key : std::true_type {}; This property is useful when the source language represents names differently in the source code and the generated code. -For example, {cpp} function names in the generated code are "mangled" in an -implementation-defined way. +For example, {cpp} function names and the names of static variables at global +scope are "mangled" in an implementation-defined way in the generated code. The precise meaning of this property is defined by each source language, but in general it allows the application to supply a list of names as they appear in the source code.