From 02ed2120396ceddab45a644401bc0e59622c959c Mon Sep 17 00:00:00 2001 From: Aaron Greig Date: Tue, 24 Sep 2024 15:38:14 +0100 Subject: [PATCH] [UR] Make multi-device compile exp functions core. --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 16 +++++----- sycl/source/backend.cpp | 19 ++---------- sycl/source/detail/kernel_bundle_impl.hpp | 9 ++---- .../program_manager/program_manager.cpp | 30 ++++--------------- .../test-e2e/ESIMD/esimd_check_vc_codegen.cpp | 7 ++--- .../test-e2e/Graph/Explicit/kernel_bundle.cpp | 2 +- .../Graph/RecordReplay/kernel_bundle.cpp | 2 +- .../helpers/RuntimeLinkingCommon.hpp | 8 ++--- .../kernel-and-program/KernelBuildOptions.cpp | 12 ++++---- .../kernel-and-program/OutOfResources.cpp | 2 +- .../program_manager/DynamicLinking.cpp | 26 ---------------- .../passing_link_and_compile_options.cpp | 20 ++++++------- 12 files changed, 44 insertions(+), 109 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 410e1a4546252..e3c231b01b4da 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,14 +116,14 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 7a2caca559da81620c7430521873045bf42eafc1 - # Merge: 1d1808a4 96f66e0f - # Author: Piotr Balcer - # Date: Tue Sep 24 13:00:46 2024 +0200 - # Merge pull request #2117 from pbalcer/fix-filter-out-same-cmdlists - # Fix urEnqueueEventsWaitWithBarrier when used with interop events - set(UNIFIED_RUNTIME_TAG 7a2caca559da81620c7430521873045bf42eafc1) + set(UNIFIED_RUNTIME_REPO "https://github.com/aarongreig/unified-runtime.git") + # commit f5c907a0f74fd6729be5c2e137144f1a43f87111 + # Merge: 9ca3ec7 be38e567 + # Author: aarongreig + # Date: Mon Sep 23 08:27:12 2024 -0700 + # Merge pull request #1830 from JackAKirk/hip-set-device + # [hip] Remove deprecated hip APIs, simplify urContext + set(UNIFIED_RUNTIME_TAG aaron/makeDeviceCompileExtCore) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") # Due to the use of dependentloadflag and no installer for UMF and hwloc we need diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp index a6e422264178b..803183d90bc4a 100644 --- a/sycl/source/backend.cpp +++ b/sycl/source/backend.cpp @@ -225,22 +225,14 @@ make_kernel_bundle(ur_native_handle_t NativeHandle, switch (BinaryType) { case (UR_PROGRAM_BINARY_TYPE_NONE): if (State == bundle_state::object) { - auto Res = Adapter->call_nocheck( + auto Res = Adapter->call_nocheck( UrProgram, 1, &Dev, nullptr); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), UrProgram, nullptr); - } Adapter->checkUrResult(Res); } else if (State == bundle_state::executable) { - auto Res = Adapter->call_nocheck( + auto Res = Adapter->call_nocheck( UrProgram, 1, &Dev, nullptr); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), UrProgram, nullptr); - } Adapter->checkUrResult(Res); } @@ -254,14 +246,9 @@ make_kernel_bundle(ur_native_handle_t NativeHandle, detail::codeToString(UR_RESULT_ERROR_INVALID_VALUE)); if (State == bundle_state::executable) { ur_program_handle_t UrLinkedProgram = nullptr; - auto Res = Adapter->call_nocheck( + auto Res = Adapter->call_nocheck( ContextImpl->getHandleRef(), 1, &Dev, 1, &UrProgram, nullptr, &UrLinkedProgram); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), 1, &UrProgram, nullptr, - &UrLinkedProgram); - } Adapter->checkUrResult(Res); if (UrLinkedProgram != nullptr) { UrProgram = UrLinkedProgram; diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 38b6bb1deb920..dcb1d8e42614e 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -462,13 +462,8 @@ class kernel_bundle_impl { "urProgramCreateWithIL resulted in a null program handle."); std::string XsFlags = extractXsFlags(BuildOptions); - auto Res = Adapter->call_nocheck( - UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str()); - } - Adapter->checkUrResult(Res); + Adapter->call(UrProgram, DeviceVec.size(), + DeviceVec.data(), XsFlags.c_str()); // Get the number of kernels in the program. size_t NumKernels; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 2fa9e75be6f0c..87319308665fd 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1158,13 +1158,8 @@ static ur_result_t doCompile(const AdapterPtr &Adapter, const char *Opts) { // Try to compile with given devices, fall back to compiling with the program // context if unsupported by the adapter - auto Result = Adapter->call_nocheck( - Program, NumDevs, Devs, Opts); - if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - return Adapter->call_nocheck(Ctx, Program, - Opts); - } - return Result; + return Adapter->call_nocheck(Program, NumDevs, + Devs, Opts); } static ur_program_handle_t loadDeviceLibFallback(const ContextImplPtr Context, @@ -1563,13 +1558,9 @@ ProgramManager::ProgramPtr ProgramManager::build( const std::string &Options = LinkOptions.empty() ? CompileOptions : (CompileOptions + " " + LinkOptions); - ur_result_t Error = Adapter->call_nocheck( + ur_result_t Error = Adapter->call_nocheck( Program.get(), /*num devices =*/1, &Device, Options.c_str()); - if (Error == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Error = Adapter->call_nocheck( - Context->getHandleRef(), Program.get(), Options.c_str()); - } if (Error != UR_RESULT_SUCCESS) throw detail::set_ur_error( @@ -1599,15 +1590,10 @@ ProgramManager::ProgramPtr ProgramManager::build( ur_program_handle_t LinkedProg = nullptr; auto doLink = [&] { - auto Res = Adapter->call_nocheck( + auto Res = Adapter->call_nocheck( Context->getHandleRef(), /*num devices =*/1, &Device, LinkPrograms.size(), LinkPrograms.data(), LinkOptions.c_str(), &LinkedProg); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - Context->getHandleRef(), LinkPrograms.size(), LinkPrograms.data(), - LinkOptions.c_str(), &LinkedProg); - } return Res; }; ur_result_t Error = doLink(); @@ -2422,16 +2408,10 @@ ProgramManager::link(const device_image_plain &DeviceImage, ur_program_handle_t LinkedProg = nullptr; auto doLink = [&] { - auto Res = Adapter->call_nocheck( + return Adapter->call_nocheck( ContextImpl->getHandleRef(), URDevices.size(), URDevices.data(), URPrograms.size(), URPrograms.data(), LinkOptionsStr.c_str(), &LinkedProg); - if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { - Res = Adapter->call_nocheck( - ContextImpl->getHandleRef(), URPrograms.size(), URPrograms.data(), - LinkOptionsStr.c_str(), &LinkedProg); - } - return Res; }; ur_result_t Error = doLink(); if (Error == UR_RESULT_ERROR_OUT_OF_RESOURCES || diff --git a/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp b/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp index 09281973850c2..cc9489c03cb70 100644 --- a/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp +++ b/sycl/test-e2e/ESIMD/esimd_check_vc_codegen.cpp @@ -91,7 +91,6 @@ int main(void) { return err_cnt > 0 ? 1 : 0; } -// Don't use -NEXT here to split the line because we need to allow for the -// possbility of a BuildExp( that fails with UNSUPPORTED followed by a Build( -// that results in SUCCESS -// CHECK: ---> urProgramBuild{{(Exp)?}}({{.*}}-vc-codegen{{.*}} -> UR_RESULT_SUCCESS +// CHECK: ---> urProgramBuild( +// CHECK-SAME: .pOptions = {{.*}}-vc-codegen +// CHECK-SAME: -> UR_RESULT_SUCCESS diff --git a/sycl/test-e2e/Graph/Explicit/kernel_bundle.cpp b/sycl/test-e2e/Graph/Explicit/kernel_bundle.cpp index a43f7bb47d37f..4c918414bc0fb 100644 --- a/sycl/test-e2e/Graph/Explicit/kernel_bundle.cpp +++ b/sycl/test-e2e/Graph/Explicit/kernel_bundle.cpp @@ -9,7 +9,7 @@ // CHECK-SAME: .phProgram = {{.*}} ([[PROGRAM_HANDLE1:[0-9a-fA-Fx]+]]) // -// CHECK:---> urProgramBuildExp( +// CHECK:---> urProgramBuild( // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]] // // CHECK:---> urProgramRetain(.hProgram = [[PROGRAM_HANDLE1]]) -> UR_RESULT_SUCCESS diff --git a/sycl/test-e2e/Graph/RecordReplay/kernel_bundle.cpp b/sycl/test-e2e/Graph/RecordReplay/kernel_bundle.cpp index d347cc58b0754..a90c3180106a3 100644 --- a/sycl/test-e2e/Graph/RecordReplay/kernel_bundle.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/kernel_bundle.cpp @@ -9,7 +9,7 @@ // CHECK-SAME: .phProgram = {{.*}} ([[PROGRAM_HANDLE1:[0-9a-fA-Fx]+]]) // CHECK-SAME: -> UR_RESULT_SUCCESS; // -// CHECK:---> urProgramBuildExp( +// CHECK:---> urProgramBuild( // CHECK-SAME: .hProgram = [[PROGRAM_HANDLE1]] // // CHECK:---> urProgramRetain( diff --git a/sycl/unittests/helpers/RuntimeLinkingCommon.hpp b/sycl/unittests/helpers/RuntimeLinkingCommon.hpp index 87f0f980119f0..775484bc45441 100644 --- a/sycl/unittests/helpers/RuntimeLinkingCommon.hpp +++ b/sycl/unittests/helpers/RuntimeLinkingCommon.hpp @@ -48,8 +48,8 @@ static ur_result_t redefined_urProgramCreateWithBinary(void *pParams) { return UR_RESULT_SUCCESS; } -static ur_result_t redefined_urProgramLinkExp(void *pParams) { - auto Params = *static_cast(pParams); +static ur_result_t redefined_urProgramLink(void *pParams) { + auto Params = *static_cast(pParams); unsigned ResProgram = 1; auto Programs = *Params.pphPrograms; for (uint32_t I = 0; I < *Params.pcount; ++I) { @@ -82,8 +82,8 @@ static void setupRuntimeLinkingMock() { redefined_urProgramCreateWithIL); mock::getCallbacks().set_replace_callback( "urProgramCreateWithBinary", redefined_urProgramCreateWithBinary); - mock::getCallbacks().set_replace_callback("urProgramLinkExp", - redefined_urProgramLinkExp); + mock::getCallbacks().set_replace_callback("urProgramLink", + redefined_urProgramLink); mock::getCallbacks().set_replace_callback("urKernelCreate", redefined_urKernelCreate); } diff --git a/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp b/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp index 2fb0d69da97b9..2b6e109dd04ef 100644 --- a/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp +++ b/sycl/unittests/kernel-and-program/KernelBuildOptions.cpp @@ -35,7 +35,7 @@ struct KernelInfo : public unittest::MockKernelInfoBase { } // namespace sycl static ur_result_t redefinedProgramBuild(void *pParams) { - auto params = *static_cast(pParams); + auto params = *static_cast(pParams); if (*params.ppOptions) BuildOpts = *params.ppOptions; else @@ -44,7 +44,7 @@ static ur_result_t redefinedProgramBuild(void *pParams) { } static ur_result_t redefinedProgramCompile(void *pParams) { - auto params = *static_cast(pParams); + auto params = *static_cast(pParams); if (*params.ppOptions) BuildOpts = *params.ppOptions; else @@ -53,7 +53,7 @@ static ur_result_t redefinedProgramCompile(void *pParams) { } static ur_result_t redefinedProgramLink(void *pParams) { - auto params = *static_cast(pParams); + auto params = *static_cast(pParams); if (*params.ppOptions) BuildOpts = *params.ppOptions; else @@ -63,11 +63,11 @@ static ur_result_t redefinedProgramLink(void *pParams) { static void setupCommonMockAPIs(sycl::unittest::UrMock<> &Mock) { using namespace sycl::detail; - mock::getCallbacks().set_before_callback("urProgramCompileExp", + mock::getCallbacks().set_before_callback("urProgramCompile", &redefinedProgramCompile); - mock::getCallbacks().set_before_callback("urProgramLinkExp", + mock::getCallbacks().set_before_callback("urProgramLink", &redefinedProgramLink); - mock::getCallbacks().set_before_callback("urProgramBuildExp", + mock::getCallbacks().set_before_callback("urProgramBuild", &redefinedProgramBuild); } diff --git a/sycl/unittests/kernel-and-program/OutOfResources.cpp b/sycl/unittests/kernel-and-program/OutOfResources.cpp index b0b6e877ebe77..5cd19bdcb9807 100644 --- a/sycl/unittests/kernel-and-program/OutOfResources.cpp +++ b/sycl/unittests/kernel-and-program/OutOfResources.cpp @@ -127,7 +127,7 @@ TEST_P(OutOfResourcesTestSuite, urProgramLink) { nProgramLink = 0; sycl::unittest::UrMock<> Mock; ErrorCode = GetParam(); - mock::getCallbacks().set_before_callback("urProgramLinkExp", + mock::getCallbacks().set_before_callback("urProgramLink", &redefinedProgramLink); sycl::platform Plt{sycl::platform()}; diff --git a/sycl/unittests/program_manager/DynamicLinking.cpp b/sycl/unittests/program_manager/DynamicLinking.cpp index 3fd569c8ef392..ea765a1128ee6 100644 --- a/sycl/unittests/program_manager/DynamicLinking.cpp +++ b/sycl/unittests/program_manager/DynamicLinking.cpp @@ -218,30 +218,4 @@ TEST(DynamicLinking, AheadOfTimeUnsupported) { } } -static ur_result_t redefined_urProgramCompileExp(void *pParams) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; -} - -TEST(DynamicLinking, UnsupportedCompileExp) { - sycl::unittest::UrMock<> Mock; - setupRuntimeLinkingMock(); - mock::getCallbacks().set_replace_callback("urProgramCompileExp", - redefined_urProgramCompileExp); - - sycl::platform Plt = sycl::platform(); - sycl::queue Q(Plt.get_devices()[0]); - - CapturedLinkingData.clear(); - - Q.single_task([=]() {}); - ASSERT_EQ(CapturedLinkingData.NumOfUrProgramCreateCalls, 3u); - // Both programs should be linked together. - ASSERT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); - ASSERT_TRUE(CapturedLinkingData.LinkedProgramsContains( - {BASIC_CASE_PRG, BASIC_CASE_PRG_DEP, BASIC_CASE_PRG_DEP_DEP})); - // And the linked program should be used to create a kernel. - ASSERT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, - BASIC_CASE_PRG * BASIC_CASE_PRG_DEP * BASIC_CASE_PRG_DEP_DEP); -} - } // anonymous namespace diff --git a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp index 893e8bdb41ee7..7e1289905bddb 100644 --- a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp +++ b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp @@ -85,7 +85,7 @@ generateEAMTestKernelImage(std::string _cmplOptions, std::string _lnkOptions) { } inline ur_result_t redefinedProgramLink(void *pParams) { - auto params = *static_cast(pParams); + auto params = *static_cast(pParams); assert(*params.ppOptions != nullptr); auto add_link_opts = std::string(*params.ppOptions); if (!add_link_opts.empty()) { @@ -97,7 +97,7 @@ inline ur_result_t redefinedProgramLink(void *pParams) { } inline ur_result_t redefinedProgramCompile(void *pParams) { - auto params = *static_cast(pParams); + auto params = *static_cast(pParams); assert(*params.ppOptions != nullptr); auto add_compile_opts = std::string(*params.ppOptions); if (!add_compile_opts.empty()) { @@ -109,7 +109,7 @@ inline ur_result_t redefinedProgramCompile(void *pParams) { } inline ur_result_t redefinedProgramBuild(void *pParams) { - auto params = *static_cast(pParams); + auto params = *static_cast(pParams); assert(*params.ppOptions != nullptr); current_build_opts = std::string(*params.ppOptions); return UR_RESULT_SUCCESS; @@ -118,9 +118,9 @@ inline ur_result_t redefinedProgramBuild(void *pParams) { TEST(Link_Compile_Options, compile_link_Options_Test_empty_options) { sycl::unittest::UrMock<> Mock; sycl::platform Plt = sycl::platform(); - mock::getCallbacks().set_before_callback("urProgramCompileExp", + mock::getCallbacks().set_before_callback("urProgramCompile", &redefinedProgramCompile); - mock::getCallbacks().set_before_callback("urProgramLinkExp", + mock::getCallbacks().set_before_callback("urProgramLink", &redefinedProgramLink); const sycl::device Dev = Plt.get_devices()[0]; current_link_options.clear(); @@ -145,9 +145,9 @@ TEST(Link_Compile_Options, compile_link_Options_Test_empty_options) { TEST(Link_Compile_Options, compile_link_Options_Test_filled_options) { sycl::unittest::UrMock<> Mock; sycl::platform Plt = sycl::platform(); - mock::getCallbacks().set_before_callback("urProgramCompileExp", + mock::getCallbacks().set_before_callback("urProgramCompile", &redefinedProgramCompile); - mock::getCallbacks().set_before_callback("urProgramLinkExp", + mock::getCallbacks().set_before_callback("urProgramLink", &redefinedProgramLink); const sycl::device Dev = Plt.get_devices()[0]; current_link_options.clear(); @@ -180,11 +180,11 @@ TEST(Link_Compile_Options, compile_link_Options_Test_filled_options) { TEST(Link_Compile_Options, check_sycl_build) { sycl::unittest::UrMock<> Mock; sycl::platform Plt = sycl::platform(); - mock::getCallbacks().set_before_callback("urProgramCompileExp", + mock::getCallbacks().set_before_callback("urProgramCompile", &redefinedProgramCompile); - mock::getCallbacks().set_before_callback("urProgramLinkExp", + mock::getCallbacks().set_before_callback("urProgramLink", &redefinedProgramLink); - mock::getCallbacks().set_before_callback("urProgramBuildExp", + mock::getCallbacks().set_before_callback("urProgramBuild", &redefinedProgramBuild); const sycl::device Dev = Plt.get_devices()[0]; current_link_options.clear();