From 750627b213578dae17064ea89513efd09790d3fb Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Mon, 17 Nov 2025 14:25:34 -0800 Subject: [PATCH] [SYCL] Remove the old mechanism of instantiating kernel on the host --- sycl/include/sycl/detail/cg_types.hpp | 72 ------------------- sycl/test/abi/vtable.cpp | 12 ---- .../basic_tests/single_task_error_message.cpp | 9 --- 3 files changed, 93 deletions(-) diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp index 3368e333ac94d..6125cf327e03c 100644 --- a/sycl/include/sycl/detail/cg_types.hpp +++ b/sycl/include/sycl/detail/cg_types.hpp @@ -152,10 +152,6 @@ class HostKernelBase { // Used to extract captured variables. virtual char *getPtr() = 0; virtual ~HostKernelBase() noexcept = default; -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES - // NOTE: InstatiateKernelOnHost() should not be called. - virtual void InstantiateKernelOnHost() = 0; -#endif }; // Class which stores specific lambda object. @@ -170,69 +166,6 @@ class HostKernel : public HostKernelBase { char *getPtr() override { return reinterpret_cast(&MKernel); } ~HostKernel() noexcept override = default; - -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES - // This function is needed for host-side compilation to keep kernels - // instantitated. This is important for debuggers to be able to associate - // kernel code instructions with source code lines. - // NOTE: InstatiateKernelOnHost() should not be called. - void InstantiateKernelOnHost() override { - using IDBuilder = sycl::detail::Builder; - constexpr bool HasKernelHandlerArg = - KernelLambdaHasKernelHandlerArgT::value; - if constexpr (std::is_same_v) { - runKernelWithoutArg(MKernel, std::bool_constant()); - } else if constexpr (std::is_same_v>) { - sycl::id ID = InitializedVal::template get<0>(); - runKernelWithArg( - MKernel, ID, std::bool_constant()); - } else if constexpr (std::is_same_v> || - std::is_same_v>) { - constexpr bool HasOffset = - std::is_same_v>; - if constexpr (!HasOffset) { - KernelArgType Item = IDBuilder::createItem( - InitializedVal::template get<1>(), - InitializedVal::template get<0>()); - runKernelWithArg( - MKernel, Item, std::bool_constant()); - } else { - KernelArgType Item = IDBuilder::createItem( - InitializedVal::template get<1>(), - InitializedVal::template get<0>(), - InitializedVal::template get<0>()); - runKernelWithArg( - MKernel, Item, std::bool_constant()); - } - } else if constexpr (std::is_same_v>) { - sycl::range Range = InitializedVal::template get<1>(); - sycl::id ID = InitializedVal::template get<0>(); - sycl::group Group = - IDBuilder::createGroup(Range, Range, Range, ID); - sycl::item GlobalItem = - IDBuilder::createItem(Range, ID, ID); - sycl::item LocalItem = - IDBuilder::createItem(Range, ID); - KernelArgType NDItem = - IDBuilder::createNDItem(GlobalItem, LocalItem, Group); - runKernelWithArg( - MKernel, NDItem, std::bool_constant()); - } else if constexpr (std::is_same_v>) { - sycl::range Range = InitializedVal::template get<1>(); - sycl::id ID = InitializedVal::template get<0>(); - KernelArgType Group = - IDBuilder::createGroup(Range, Range, Range, ID); - runKernelWithArg( - MKernel, Group, std::bool_constant()); - } else { - // Assume that anything else can be default-constructed. If not, this - // should fail to compile and the implementor should implement a generic - // case for the new argument type. - runKernelWithArg( - MKernel, KernelArgType{}, std::bool_constant()); - } - } -#endif }; // the class keeps reference to a lambda allocated externally on stack @@ -243,11 +176,6 @@ class HostKernelRefBase : public HostKernelBase { HostKernelRefBase &operator=(const HostKernelRefBase &) = delete; virtual std::unique_ptr takeOrCopyOwnership() const = 0; -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES - // The kernels that are passed via HostKernelRefBase are instantiated along - // ctor call with GetInstantiateKernelOnHostPtr(). - void InstantiateKernelOnHost() override {} -#endif }; // Primary template for movable objects. diff --git a/sycl/test/abi/vtable.cpp b/sycl/test/abi/vtable.cpp index bdfcd4dc5f16c..f1ff7e02f1900 100644 --- a/sycl/test/abi/vtable.cpp +++ b/sycl/test/abi/vtable.cpp @@ -8,18 +8,6 @@ // Changing vtable breaks ABI. If this test fails, please, refer to ABI Policy // Guide for further instructions. -void foo(sycl::detail::HostKernelBase &HKB) { - HKB.InstantiateKernelOnHost(); -} -// CHECK: Vtable for 'sycl::detail::HostKernelBase' (6 entries). -// CHECK-NEXT: 0 | offset_to_top (0) -// CHECK-NEXT: 1 | sycl::detail::HostKernelBase RTTI -// CHECK-NEXT: -- (sycl::detail::HostKernelBase, 0) vtable address -- -// CHECK-NEXT: 2 | char *sycl::detail::HostKernelBase::getPtr() [pure] -// CHECK-NEXT: 3 | sycl::detail::HostKernelBase::~HostKernelBase() [complete] -// CHECK-NEXT: 4 | sycl::detail::HostKernelBase::~HostKernelBase() [deleting] -// CHECK-NEXT: 5 | void sycl::detail::HostKernelBase::InstantiateKernelOnHost() [pure] - void foo(sycl::detail::PropertyWithDataBase *Prop) { delete Prop; } // CHECK: Vtable for 'sycl::detail::PropertyWithDataBase' (4 entries). // CHECK-NEXT: 0 | offset_to_top (0) diff --git a/sycl/test/basic_tests/single_task_error_message.cpp b/sycl/test/basic_tests/single_task_error_message.cpp index 7c6f020741247..7389296645d63 100644 --- a/sycl/test/basic_tests/single_task_error_message.cpp +++ b/sycl/test/basic_tests/single_task_error_message.cpp @@ -12,9 +12,6 @@ int main() { .single_task([&](sycl::handler &cgh) { // expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}} // expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}} - // TODO Investigate why this function template is not instantiated - // (if this is expected). - // expected-error@sycl/detail/cg_types.hpp:* {{no matching function for call to 'runKernelWithoutArg'}} }) .wait(); } @@ -31,9 +28,6 @@ int main() { [&](sycl::handler &cgh) { // expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}} // expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}} - // TODO Investigate why this function template is not - // instantiated (if this is expected). - // expected-error@sycl/detail/cg_types.hpp:* {{no matching function for call to 'runKernelWithoutArg'}} }) .wait(); } @@ -50,9 +44,6 @@ int main() { [&](sycl::handler &cgh) { // expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}} // expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}} - // TODO Investigate why this function template is not - // instantiated (if this is expected). - // expected-error@sycl/detail/cg_types.hpp:* {{no matching function for call to 'runKernelWithoutArg'}} }) .wait(); }