diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp index 6e5b933fa48d7..fcde2d334d3b8 100644 --- a/sycl/source/detail/kernel_impl.hpp +++ b/sycl/source/detail/kernel_impl.hpp @@ -115,6 +115,9 @@ class kernel_impl { typename Param::return_type get_info(const device &Device, const range<3> &WGSize) const; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // This function is unused and should be removed in the next ABI breaking. + /// Query queue/launch-specific information from a kernel using the /// info::kernel_queue_specific descriptor for a specific Queue. /// @@ -122,6 +125,7 @@ class kernel_impl { /// \return depends on information being queried. template typename Param::return_type ext_oneapi_get_info(queue Queue) const; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES /// Query queue/launch-specific information from a kernel using the /// info::kernel_queue_specific descriptor for a specific Queue and values. @@ -440,6 +444,9 @@ inline typename ext::intel::info::kernel_device_specific::spill_memory_size:: getAdapter()); } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +// These functions are unused and should be removed in the next ABI breaking. + template <> inline typename syclex::info::kernel_queue_specific::max_work_group_size:: return_type @@ -491,6 +498,8 @@ ADD_TEMPLATE_METHOD_SPEC(3) #undef ADD_TEMPLATE_METHOD_SPEC +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + #define ADD_TEMPLATE_METHOD_SPEC(QueueSpec, Num, Kind, Reg) \ template <> \ inline typename syclex::info::kernel_queue_specific::QueueSpec::return_type \ diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 724dcc2956734..2d87758602f53 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -303,6 +303,9 @@ bool Command::isHostTask() const { CGType::CodeplayHostTask); } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +// This function is unused and should be removed in the next ABI-breaking +// window. bool Command::isFusable() const { if ((MType != CommandType::RUN_CG)) { return false; @@ -312,6 +315,7 @@ bool Command::isFusable() const { (!static_cast(CG).MKernelIsCooperative) && (!static_cast(CG).MKernelUsesClusterLaunch); } +#endif // __INTEL_PREVIEW_BREAKING_CHANGES static void flushCrossQueueDeps(const std::vector &EventImpls, const QueueImplPtr &Queue) { diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 742557d1b65ba..b8563e7596369 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -251,7 +251,11 @@ class Command { bool isHostTask() const; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // This function is unused and should be removed in the next ABI-breaking + // window. bool isFusable() const; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES protected: QueueImplPtr MQueue; diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp index d01d6d9e765c0..8d743b643feb5 100644 --- a/sycl/source/kernel.cpp +++ b/sycl/source/kernel.cpp @@ -114,11 +114,15 @@ template __SYCL_EXPORT uint32_t kernel::get_info( const device &, const sycl::range<3> &) const; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +// This function is unused and should be removed in the next ABI-breaking +// window. template typename detail::is_kernel_queue_specific_info_desc::return_type kernel::ext_oneapi_get_info(queue Queue) const { return impl->ext_oneapi_get_info(std::move(Queue)); } +#endif // __INTEL_PREVIEW_BREAKING_CHANGES template typename detail::is_kernel_queue_specific_info_desc::return_type @@ -162,6 +166,10 @@ kernel::ext_oneapi_get_info(queue Queue, const range<3> &WorkGroupSize, DynamicLocalMemorySize); } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +// These functions are unused and should be removed in the next ABI-breaking +// window. + template __SYCL_EXPORT typename ext::oneapi::experimental::info:: kernel_queue_specific::max_work_group_size::return_type kernel::ext_oneapi_get_info>(queue Queue) const; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + template __SYCL_EXPORT typename ext::oneapi::experimental::info:: kernel_queue_specific::max_sub_group_size::return_type kernel::ext_oneapi_get_info +#include + +// Include helpers for device image, kernel info, and Unified Runtime (UR) mocks +#include "helpers/MockDeviceImage.hpp" +#include "helpers/MockKernelInfo.hpp" +#include "helpers/UrMock.hpp" + +// Define a mock kernel class with several operator() overloads for different +// SYCL item types +class QueryKernel { +public: + void operator()() const {} + void operator()(sycl::item<1>) const {} + void operator()(sycl::nd_item<1> Item) const {} +}; + +// Specialize KernelInfo for QueryKernel to provide mock metadata for the kernel +namespace sycl { +inline namespace _V1 { +namespace detail { +template <> +struct KernelInfo : public unittest::MockKernelInfoBase { + static constexpr const char *getName() { return "QueryKernel"; } + static constexpr int64_t getKernelSize() { return sizeof(QueryKernel); } + static constexpr const char *getFileName() { return "QueryKernel.hpp"; } + static constexpr const char *getFunctionName() { + return "QueryKernelFunctionName"; + } + static constexpr unsigned getLineNumber() { return 1; } + static constexpr unsigned getColumnNumber() { return 1; } +}; +} // namespace detail +} // namespace _V1 +} // namespace sycl + +// Test that querying max_num_work_groups with an invalid (zero) work-group size +// throws the correct exception +TEST(RootGroupTests, InvalidWorkGroupSize) { + namespace syclex = sycl::ext::oneapi::experimental; + + // Create a mock device image containing the QueryKernel + sycl::unittest::MockDeviceImage Img = + sycl::unittest::generateDefaultImage({"QueryKernel"}); + const sycl::unittest::MockDeviceImageArray<1> ImgArray{&Img}; + const sycl::unittest::UrMock<> Mock; + + const sycl::queue q; + // Get the kernel bundle and kernel object for QueryKernel + const auto bundle = + sycl::get_kernel_bundle(q.get_context()); + const auto kernel = bundle.get_kernel(); + try { + // Attempt to query max_num_work_groups with a zero work-group size + kernel.ext_oneapi_get_info< + syclex::info::kernel_queue_specific::max_num_work_groups>(q, {0}, 0); + FAIL() << "The ext_oneapi_get_info query should have thrown."; + } catch (const sycl::exception &e) { + // Check that the correct error code and message are returned + EXPECT_EQ(e.code(), sycl::make_error_code(sycl::errc::invalid)); + EXPECT_STREQ(e.what(), "The launch work-group size cannot be zero."); + } +} + +// Test that querying max_num_work_groups with a valid work-group size returns +// the expected value +TEST(RootGroupTests, ValidNumWorkGroupsQuery) { + namespace syclex = sycl::ext::oneapi::experimental; + + // Create a mock device image containing the QueryKernel + sycl::unittest::MockDeviceImage Img = + sycl::unittest::generateDefaultImage({"QueryKernel"}); + const sycl::unittest::MockDeviceImageArray<1> ImgArray{&Img}; + const sycl::unittest::UrMock<> Mock; + + // Set up a mock callback to return a specific group count when queried + constexpr std::size_t mock_group_count = 42; + mock::getCallbacks().set_replace_callback( + "urKernelSuggestMaxCooperativeGroupCountExp", [](void *pParams) { + auto params = static_cast< + ur_kernel_suggest_max_cooperative_group_count_exp_params_t *>( + pParams); + **params->ppGroupCountRet = mock_group_count; + return UR_RESULT_SUCCESS; + }); + + const sycl::queue q; + // Get the kernel bundle and kernel object for QueryKernel + const auto bundle = + sycl::get_kernel_bundle(q.get_context()); + const auto kernel = bundle.get_kernel(); + // Query max_num_work_groups with a valid work-group size (1) + const auto maxWGs = kernel.ext_oneapi_get_info< + syclex::info::kernel_queue_specific::max_num_work_groups>(q, {1}, 0); + // Check that the returned value matches the mock group count + EXPECT_EQ(maxWGs, mock_group_count); +}