diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index 8fb38aa01edfe..150206625f586 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -86,6 +86,7 @@ def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group" def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">; def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">; def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">; +def AspectExt_intel_spill_memory_size : Aspect<"ext_intel_spill_memory_size">; // Deprecated aspects def AspectInt64_base_atomics : Aspect<"int64_base_atomics">; def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">; @@ -150,7 +151,8 @@ def : TargetInfo<"__TestAspectList", AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group, AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_atomic16, - AspectExt_oneapi_virtual_functions], + AspectExt_oneapi_virtual_functions, + AspectExt_intel_spill_memory_size], []>; // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT // match. diff --git a/sycl/doc/extensions/supported/sycl_ext_intel_kernel_queries.asciidoc b/sycl/doc/extensions/supported/sycl_ext_intel_kernel_queries.asciidoc new file mode 100644 index 0000000000000..b030436af5322 --- /dev/null +++ b/sycl/doc/extensions/supported/sycl_ext_intel_kernel_queries.asciidoc @@ -0,0 +1,139 @@ += sycl_ext_intel_kernel_queries + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en +:dpcpp: pass:[DPC++] +:endnote: —{nbsp}end{nbsp}note + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + + +== Notice + +[%hardbreaks] +Copyright (C) 2025 Intel Corporation. All rights reserved. + +Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks +of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by +permission by Khronos. + + +== Contact + +To report problems with this extension, please open a new issue at: + +https://github.com/intel/llvm/issues + + +== Dependencies + +This extension is written against the SYCL 2020 revision 9 specification. +All references below to the "core SYCL specification" or to section numbers in +the SYCL specification refer to that revision. + + +== Status + +This extension is implemented and fully supported by {dpcpp}. + + +== Overview + +This extension contains a collection of queries that provide low-level +information about kernels. +These queries generally forward directly to the backend and expose concepts that +are specific to a particular implementation. +As a result, these queries may not be supported for all devices. +Each query has an associate device aspect, which tells whether the query is +supported on that device. + + +== Specification + +=== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification. +An implementation supporting this extension must predefine the macro +`SYCL_EXT_INTEL_KERNEL_QUERIES` to one of the values defined in the table below. +Applications can test for the existence of this macro to determine if the +implementation supports this feature, or applications can test the macro's value +to determine which of the extension's features the implementation supports. + +[%header,cols="1,5"] +|=== +|Value +|Description + +|1 +|Initial version of this extension. +|=== + +=== Spill memory size + +This query returns the kernel's spill memory size that is allocated by the +compiler, as reported by Level Zero. + +==== New device aspect + +This extension adds the following new device aspect. + +[source,c++] +---- +namespace sycl { + +enum class aspect { + ext_intel_spill_memory_size + + // ... +}; + +} +---- + +''' + +`*ext_intel_spill_memory_size*` + +Indicates that the `spill_memory_size` kernel information descriptor may be used +to query kernels for this device. + +''' + +==== New device specific kernel information descriptor + +This extension adds the following information descriptor that can be used with +`kernel::get_info(const device&)`. + +''' + +`*ext::intel::info::kernel_device_specific::spill_memory_size*` + +[source,c++] +---- +namespace sycl::ext::intel::info::kernel_device_specific { +struct spill_memory_size { + using return_type = size_t; +}; +} // namespace sycl::ext::intel::info::kernel_device_specific +---- + +_Remarks:_ Template parameter to `kernel::get_info(const device&)`. + +_Returns:_ The spill memory size that is allocated by the compiler for this +kernel for the given device. + +_Throws:_ An `exception` with the `errc::feature_not_supported` error code if +the device does not have `aspect::ext_intel_spill_memory_size`. + +''' diff --git a/sycl/include/sycl/detail/info_desc_helpers.hpp b/sycl/include/sycl/detail/info_desc_helpers.hpp index 2644c98cbb083..d3310ea7978f9 100644 --- a/sycl/include/sycl/detail/info_desc_helpers.hpp +++ b/sycl/include/sycl/detail/info_desc_helpers.hpp @@ -88,6 +88,15 @@ struct IsKernelInfo #include #undef __SYCL_PARAM_TRAITS_SPEC +#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode) \ + template <> \ + struct is_##DescType##_info_desc \ + : std::true_type { \ + using return_type = Namespace::info::DescType::Desc::return_type; \ + }; +#include +#undef __SYCL_PARAM_TRAITS_SPEC + #define __SYCL_PARAM_TRAITS_SPEC(DescType, Desc, ReturnT, UrCode) \ template <> \ struct is_##DescType##_info_desc : std::true_type { \ diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 8a931dde35a71..d039e4bc2fee5 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/sycl/include/sycl/info/aspects.def @@ -72,3 +72,4 @@ __SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 78) __SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 79) __SYCL_ASPECT(ext_oneapi_atomic16, 80) __SYCL_ASPECT(ext_oneapi_virtual_functions, 81) +__SYCL_ASPECT(ext_intel_spill_memory_size, 82) diff --git a/sycl/include/sycl/info/ext_intel_kernel_info_traits.def b/sycl/include/sycl/info/ext_intel_kernel_info_traits.def new file mode 100644 index 0000000000000..4a0d7d27d43f5 --- /dev/null +++ b/sycl/include/sycl/info/ext_intel_kernel_info_traits.def @@ -0,0 +1 @@ +__SYCL_PARAM_TRAITS_SPEC(ext::intel, kernel_device_specific, spill_memory_size, size_t, UR_KERNEL_INFO_SPILL_MEM_SIZE) diff --git a/sycl/include/sycl/info/info_desc.hpp b/sycl/include/sycl/info/info_desc.hpp index aa7baa39e2f62..8fad057283ecf 100644 --- a/sycl/include/sycl/info/info_desc.hpp +++ b/sycl/include/sycl/info/info_desc.hpp @@ -251,6 +251,7 @@ struct work_item_progress_capabilities; } // namespace ext::oneapi::experimental::info::device #include #include +#include #include #include diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 5cb7fa1e29585..9d630d20dbfeb 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -786,6 +786,11 @@ bool device_impl::has(aspect Aspect) const { BE == sycl::backend::opencl; return (is_cpu() || is_gpu()) && isCompatibleBE; } + case aspect::ext_intel_spill_memory_size: { + backend BE = getBackend(); + bool isCompatibleBE = BE == sycl::backend::ext_oneapi_level_zero; + return is_gpu() && isCompatibleBE; + } } return false; // This device aspect has not been implemented yet. diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp index 1b07d866dcc4c..8f1d8d702e2b5 100644 --- a/sycl/source/detail/kernel_impl.hpp +++ b/sycl/source/detail/kernel_impl.hpp @@ -383,6 +383,23 @@ inline typename syclex::info::kernel_queue_specific::max_num_work_groups:: DynamicLocalMemorySize); } +template <> +inline typename ext::intel::info::kernel_device_specific::spill_memory_size:: + return_type + kernel_impl::get_info< + ext::intel::info::kernel_device_specific::spill_memory_size>( + const device &Device) const { + if (!Device.has(aspect::ext_intel_spill_memory_size)) + throw exception( + make_error_code(errc::feature_not_supported), + "This device does not have the ext_intel_spill_memory_size aspect"); + + return get_kernel_device_specific_info< + ext::intel::info::kernel_device_specific::spill_memory_size>( + this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), + getAdapter()); +} + } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/kernel_info.hpp b/sycl/source/detail/kernel_info.hpp index 4ec11d32fd4ac..9c550e928e723 100644 --- a/sycl/source/detail/kernel_info.hpp +++ b/sycl/source/detail/kernel_info.hpp @@ -146,6 +146,54 @@ uint32_t get_kernel_device_specific_info_with_input(ur_kernel_handle_t Kernel, return Result; } + +template <> +inline ext::intel::info::kernel_device_specific::spill_memory_size::return_type +get_kernel_device_specific_info< + ext::intel::info::kernel_device_specific::spill_memory_size>( + ur_kernel_handle_t Kernel, ur_device_handle_t Device, + const AdapterPtr &Adapter) { + size_t ResultSize = 0; + ur_kernel_info_t PropName = UR_KERNEL_INFO_SPILL_MEM_SIZE; + + // First call to get the number of device images + Adapter->call(Kernel, PropName, 0, nullptr, + &ResultSize); + + size_t DeviceCount = ResultSize / sizeof(uint32_t); + + // Second call to retrieve the data + std::vector Result(DeviceCount); + Adapter->call(Kernel, PropName, ResultSize, + Result.data(), nullptr); + + ur_program_handle_t Program; + Adapter->call(Kernel, UR_KERNEL_INFO_PROGRAM, + sizeof(ur_program_handle_t), + &Program, nullptr); + // Retrieve the associated device list + size_t URDevicesSize = 0; + Adapter->call(Program, UR_PROGRAM_INFO_DEVICES, + 0, nullptr, &URDevicesSize); + + std::vector URDevices(URDevicesSize / + sizeof(ur_device_handle_t)); + Adapter->call(Program, UR_PROGRAM_INFO_DEVICES, + URDevicesSize, URDevices.data(), + nullptr); + assert(Result.size() == URDevices.size()); + + // Map the result back to the program devices + for (size_t idx = 0; idx < URDevices.size(); ++idx) { + if (URDevices[idx] == Device) + return size_t{Result[idx]}; + } + throw exception( + make_error_code(errc::runtime), + "ext::intel::info::kernel::spill_memory_size failed to retrieve " + "the requested value"); +} + } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/ur_info_code.hpp b/sycl/source/detail/ur_info_code.hpp index 863515ec469da..9d32c60660d6b 100644 --- a/sycl/source/detail/ur_info_code.hpp +++ b/sycl/source/detail/ur_info_code.hpp @@ -72,6 +72,15 @@ template struct UrInfoCode; #include #undef __SYCL_PARAM_TRAITS_SPEC +#define __SYCL_PARAM_TRAITS_SPEC(Namespace, DescType, Desc, ReturnT, UrCode) \ + template <> struct UrInfoCode { \ + static constexpr ur_kernel_info_t value = \ + static_cast(UrCode); \ + }; + +#include +#undef __SYCL_PARAM_TRAITS_SPEC + } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/test-e2e/Basic/aspects.cpp b/sycl/test-e2e/Basic/aspects.cpp index ea1bbec27762d..eaa33c0073397 100644 --- a/sycl/test-e2e/Basic/aspects.cpp +++ b/sycl/test-e2e/Basic/aspects.cpp @@ -90,6 +90,9 @@ int main() { if (plt.has(aspect::ext_oneapi_virtual_functions)) { std::cout << " ext_oneapi_virtual_functions" << std::endl; } + if (plt.has(aspect::ext_intel_spill_memory_size)) { + std::cout << " ext_intel_spill_memory_size" << std::endl; + } } std::cout << "Passed." << std::endl; return 0; diff --git a/sycl/test-e2e/Basic/kernel_info.cpp b/sycl/test-e2e/Basic/kernel_info.cpp index f0fdad910c658..aa9aabc59c551 100644 --- a/sycl/test-e2e/Basic/kernel_info.cpp +++ b/sycl/test-e2e/Basic/kernel_info.cpp @@ -101,6 +101,12 @@ int main() { krn.get_info(dev); assert(compileNumSg <= maxNumSg); + if (dev.has(aspect::ext_intel_spill_memory_size)) { + const size_t spillMemSz = krn.get_info< + ext::intel::info::kernel_device_specific::spill_memory_size>(dev); + assert(spillMemSz >= 0); + } + // Use ext_oneapi_get_kernel_info extension and check that answers match. const size_t wgSizeExt = syclex::get_kernel_info< SingleTask, info::kernel_device_specific::work_group_size>(ctx, dev); @@ -125,6 +131,13 @@ int main() { dev); assert(compileNumSgExt == compileNumSg); + if (dev.has(aspect::ext_intel_spill_memory_size)) { + const size_t spillMemSizeExt = syclex::get_kernel_info< + SingleTask, + ext::intel::info::kernel_device_specific::spill_memory_size>(ctx, dev); + assert(spillMemSizeExt == spillMemSz); + } + // Use ext_oneapi_get_kernel_info extension with queue parameter and check the // result. const size_t wgSizeExtQ =