Skip to content

[UR] Pull in change to make multi device compile extension core. #12536

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 6 commits into
base: sycl
Choose a base branch
from
Open

[UR] Pull in change to make multi device compile extension core. #12536

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
d5cd646
[UR] Make multi-device compile experimental feature core.
aarongreig Apr 10, 2025
e2f6479
Replace exp functions on the sycl side
aarongreig Apr 10, 2025
ada2d9d
Re-generate spec and fix l0 v2 test.
aarongreig Apr 10, 2025
3d8d95d
Merge branch 'sycl' into aaron/makeMultiDeviceCompileCore
aarongreig Apr 11, 2025
c80b35f
Fix missed conflict.
aarongreig Apr 11, 2025
7b5395d
Fix unused variable.
aarongreig Apr 11, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 5 additions & 21 deletions sycl/source/backend.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -226,23 +226,13 @@ 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<UrApiKind::urProgramCompileExp>(
UrProgram, 1, &Dev, nullptr);
if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
Res = Adapter->call_nocheck<UrApiKind::urProgramCompile>(
ContextImpl->getHandleRef(), UrProgram, nullptr);
}
Adapter->checkUrResult<errc::build>(Res);
Adapter->call<errc::build, UrApiKind::urProgramCompile>(UrProgram, 1,
&Dev, nullptr);
}

else if (State == bundle_state::executable) {
auto Res = Adapter->call_nocheck<UrApiKind::urProgramBuildExp>(
UrProgram, 1, &Dev, nullptr);
if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
Res = Adapter->call_nocheck<UrApiKind::urProgramBuild>(
ContextImpl->getHandleRef(), UrProgram, nullptr);
}
Adapter->checkUrResult<errc::build>(Res);
Adapter->call<errc::build, UrApiKind::urProgramBuild>(UrProgram, 1,
&Dev, nullptr);
}

break;
Expand All @@ -255,15 +245,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<UrApiKind::urProgramLinkExp>(
Adapter->call<errc::build, UrApiKind::urProgramLink>(
ContextImpl->getHandleRef(), 1, &Dev, 1, &UrProgram, nullptr,
&UrLinkedProgram);
if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
Res = Adapter->call_nocheck<UrApiKind::urProgramLink>(
ContextImpl->getHandleRef(), 1, &UrProgram, nullptr,
&UrLinkedProgram);
}
Adapter->checkUrResult<errc::build>(Res);
if (UrLinkedProgram != nullptr) {
UrProgram = UrLinkedProgram;
}
Expand Down
7 changes: 1 addition & 6 deletions sycl/source/detail/device_image_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -944,13 +944,8 @@ class device_image_impl {
} // if(!FetchedFromCache)

std::string XsFlags = extractXsFlags(BuildOptions);
auto Res = Adapter->call_nocheck<UrApiKind::urProgramBuildExp>(
Adapter->call<errc::build, UrApiKind::urProgramBuild>(
UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str());
if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
Res = Adapter->call_nocheck<UrApiKind::urProgramBuild>(
ContextImpl->getHandleRef(), UrProgram, XsFlags.c_str());
}
Adapter->checkUrResult<errc::build>(Res);

// Get the number of kernels in the program.
size_t NumKernels;
Expand Down
53 changes: 13 additions & 40 deletions sycl/source/detail/program_manager/program_manager.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1321,17 +1321,9 @@ static const char *getDeviceLibExtensionStr(DeviceLibExt Extension) {

static ur_result_t doCompile(const AdapterPtr &Adapter,
ur_program_handle_t Program, uint32_t NumDevs,
ur_device_handle_t *Devs, ur_context_handle_t Ctx,
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<UrApiKind::urProgramCompileExp>(
Program, NumDevs, Devs, Opts);
if (Result == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
return Adapter->call_nocheck<UrApiKind::urProgramCompile>(Ctx, Program,
Opts);
}
return Result;
ur_device_handle_t *Devs, const char *Opts) {
return Adapter->call_nocheck<UrApiKind::urProgramCompile>(Program, NumDevs,
Devs, Opts);
}

static ur_program_handle_t
Expand Down Expand Up @@ -1402,9 +1394,8 @@ loadDeviceLibFallback(const ContextImplPtr Context, DeviceLibExt Extension,
// Do not use compile options for library programs: it is not clear if user
// options (image options) are supposed to be applied to library program as
// well, and what actually happens to a SPIR-V program if we apply them.
ur_result_t Error =
doCompile(Adapter, URProgram, DevicesToCompile.size(),
DevicesToCompile.data(), Context->getHandleRef(), "");
ur_result_t Error = doCompile(Adapter, URProgram, DevicesToCompile.size(),
DevicesToCompile.data(), "");
if (Error != UR_RESULT_SUCCESS) {
EraseProgramForDevices();
throw detail::set_ur_error(
Expand Down Expand Up @@ -1745,12 +1736,8 @@ ProgramManager::ProgramPtr ProgramManager::build(
const std::string &Options = LinkOptions.empty()
? CompileOptions
: (CompileOptions + " " + LinkOptions);
ur_result_t Error = Adapter->call_nocheck<UrApiKind::urProgramBuildExp>(
ur_result_t Error = Adapter->call_nocheck<UrApiKind::urProgramBuild>(
Program.get(), Devices.size(), Devices.data(), Options.c_str());
if (Error == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
Error = Adapter->call_nocheck<UrApiKind::urProgramBuild>(
Context->getHandleRef(), Program.get(), Options.c_str());
}

if (Error != UR_RESULT_SUCCESS)
throw detail::set_ur_error(
Expand All @@ -1764,32 +1751,26 @@ ProgramManager::ProgramPtr ProgramManager::build(
// Include the main program and compile/link everything together
if (!CreatedFromBinary) {
auto Res = doCompile(Adapter, Program.get(), Devices.size(), Devices.data(),
Context->getHandleRef(), CompileOptions.c_str());
CompileOptions.c_str());
Adapter->checkUrResult<errc::build>(Res);
}
LinkPrograms.push_back(Program.get());

for (ur_program_handle_t Prg : ExtraProgramsToLink) {
if (!CreatedFromBinary) {
auto Res = doCompile(Adapter, Prg, Devices.size(), Devices.data(),
Context->getHandleRef(), CompileOptions.c_str());
CompileOptions.c_str());
Adapter->checkUrResult<errc::build>(Res);
}
LinkPrograms.push_back(Prg);
}

ur_program_handle_t LinkedProg = nullptr;
auto doLink = [&] {
auto Res = Adapter->call_nocheck<UrApiKind::urProgramLinkExp>(
return Adapter->call_nocheck<UrApiKind::urProgramLink>(
Context->getHandleRef(), Devices.size(), Devices.data(),
LinkPrograms.size(), LinkPrograms.data(), LinkOptions.c_str(),
&LinkedProg);
if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
Res = Adapter->call_nocheck<UrApiKind::urProgramLink>(
Context->getHandleRef(), LinkPrograms.size(), LinkPrograms.data(),
LinkOptions.c_str(), &LinkedProg);
}
return Res;
};
ur_result_t Error = doLink();
if (Error == UR_RESULT_ERROR_OUT_OF_RESOURCES ||
Expand Down Expand Up @@ -2906,11 +2887,9 @@ ProgramManager::compile(const DevImgPlainWithDeps &ImgWithDeps,
CompileOptions, *(InputImpl->get_bin_image_ref()), Devs, Adapter);
// Should always come last!
appendCompileEnvironmentVariablesThatAppend(CompileOptions);
ur_result_t Error = doCompile(
Adapter, ObjectImpl->get_ur_program_ref(), Devs.size(),
URDevices.data(),
getSyclObjImpl(InputImpl->get_context()).get()->getHandleRef(),
CompileOptions.c_str());
ur_result_t Error =
doCompile(Adapter, ObjectImpl->get_ur_program_ref(), Devs.size(),
URDevices.data(), CompileOptions.c_str());
if (Error != UR_RESULT_SUCCESS)
throw sycl::exception(
make_error_code(errc::build),
Expand Down Expand Up @@ -3004,16 +2983,10 @@ ProgramManager::link(const DevImgPlainWithDeps &ImgWithDeps,

ur_program_handle_t LinkedProg = nullptr;
auto doLink = [&] {
auto Res = Adapter->call_nocheck<UrApiKind::urProgramLinkExp>(
return Adapter->call_nocheck<UrApiKind::urProgramLink>(
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<UrApiKind::urProgramLink>(
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 ||
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,7 +92,7 @@ int main(int argc, char **argv) {
}

// --- Check that only two JIT compilation happened:
// CHECK-COUNT-2: <--- urProgramBuildExp
// CHECK-NOT: <--- urProgramBuildExp
// CHECK-COUNT-2: <--- urProgramBuild
// CHECK-NOT: <--- urProgramBuild
// --- Check that the test completed with expected results:
// CHECK: passed
7 changes: 2 additions & 5 deletions sycl/test-e2e/ESIMD/sycl_esimd_mix.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -119,10 +119,7 @@ int main(void) {

// Regular SYCL kernel is compiled without -vc-codegen option

// Some backends will call urProgramBuild and some will call
// urProgramBuildExp depending on urProgramBuildExp support.

// CHECK-LABEL: <--- urProgramBuild{{(Exp)?}}
// CHECK-LABEL: <--- urProgramBuild
// CHECK-NOT: -vc-codegen
// CHECK-WITH-VAR: -g
// CHECK-NOT: -vc-codegen
Expand All @@ -134,7 +131,7 @@ int main(void) {
// For ESIMD kernels, -vc-codegen option is always preserved,
// regardless of SYCL_PROGRAM_COMPILE_OPTIONS value.

// CHECK-LABEL: <--- urProgramBuild{{(Exp)?}}
// CHECK-LABEL: <--- urProgramBuild
// CHECK-NO-VAR: -vc-codegen
// CHECK-WITH-VAR: -g -vc-codegen
// CHECK: {{.*}}-> UR_RESULT_SUCCESS
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Graph/Explicit/kernel_bundle.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Graph/RecordReplay/kernel_bundle.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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(
Expand Down
7 changes: 2 additions & 5 deletions sycl/test-e2e/KernelAndProgram/cache_env_vars.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,15 +22,12 @@
// CPU OCL JIT 0.12 0.12 0.16 1.1 16
// CPU OCL Cache 0.01 0.01 0.01 0.02 0.08

// Some backends will call urProgramBuild and some will call
// urProgramBuildExp depending on urProgramBuildExp support.

// CHECK-BUILD-NOT: <--- urProgramCreateWithBinary(
// CHECK-BUILD: <--- urProgramCreateWithIL(
// CHECK-BUILD: <--- urProgramBuild{{(Exp)?}}(
// CHECK-BUILD: <--- urProgramBuild(

// CHECK-CACHE-NOT: <--- urProgramCreateWithIL(
// CHECK-CACHE: <--- urProgramCreateWithBinary(
// CHECK-CACHE: <--- urProgramBuild{{(Exp)?}}(
// CHECK-CACHE: <--- urProgramBuild(

#include "cache_env_vars.hpp"
6 changes: 2 additions & 4 deletions sycl/test-e2e/KernelAndProgram/cache_env_vars_lin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,14 +22,12 @@
// RUN: env SYCL_CACHE_PERSISTENT=1 HOME=%t/cache_dir SYCL_UR_TRACE=2 env -u XDG_CACHE_HOME env -u SYCL_CACHE_DIR %{run} %t.out | FileCheck %s --check-prefixes=CHECK-BUILD
// RUN: env SYCL_CACHE_PERSISTENT=1 HOME=%t/cache_dir SYCL_UR_TRACE=2 env -u XDG_CACHE_HOME env -u SYCL_CACHE_DIR %{run} %t.out | FileCheck %s --check-prefixes=CHECK-CACHE

// Some backends will call urProgramBuild and some will call urProgramBuildExp depending on urProgramBuildExp support.

// CHECK-BUILD-NOT: <--- urProgramCreateWithBinary(
// CHECK-BUILD: <--- urProgramCreateWithIL(
// CHECK-BUILD: <--- urProgramBuild{{(Exp)?}}(
// CHECK-BUILD: <--- urProgramBuild(

// CHECK-CACHE-NOT: <--- urProgramCreateWithIL(
// CHECK-CACHE: <--- urProgramCreateWithBinary(
// CHECK-CACHE: <--- urProgramBuild{{(Exp)?}}(
// CHECK-CACHE: <--- urProgramBuild(

#include "cache_env_vars.hpp"
10 changes: 3 additions & 7 deletions sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,19 +3,15 @@
// RUN: env SYCL_UR_TRACE=2 %{run} %t.out | FileCheck %s
// UNSUPPORTED: hip

// Note that the UR call might be urProgramBuild OR urProgramBuildExp .
// The same is true for Compile and Link.
// We want the first match. Don't put parentheses after.

#include "kernel-bundle-merge-options.hpp"

// CHECK: <--- urProgramBuild
// CHECK: <--- urProgramBuild(
// CHECK-SAME: -g

// CHECK: <--- urProgramCompile
// CHECK: <--- urProgramCompile(
// CHECK-SAME: -g

// TODO: Uncomment when build options are properly passed to link
// commands for kernel_bundle
// xCHECK: <--- urProgramLink
// xCHECK: <--- urProgramLink(
// xCHECK-SAME: -g
2 changes: 1 addition & 1 deletion sycl/test-e2e/KernelAndProgram/target_register_alloc_mode.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,7 @@
// RUN: env SYCL_UR_TRACE=2 %{run} %t_without.out 2>&1 | FileCheck %if system-windows %{ --implicit-check-not=-ze-intel-enable-auto-large-GRF-mode %} %else %{ --check-prefix=CHECK-OPT %} %s
// RUN: env SYCL_UR_TRACE=2 %{run} %t_default.out 2>&1 | FileCheck --implicit-check-not=-ze-intel-enable-auto-large-GRF-mode %s

// CHECK-OPT: <--- urProgramBuildExp(
// CHECK-OPT: <--- urProgramBuild(
// CHECK-SAME-OPT: -ze-intel-enable-auto-large-GRF-mode

#include <sycl/detail/core.hpp>
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/ProgramManager/multi_device_bundle/build_twice.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,10 +33,10 @@ int main() {
sycl::kernel_bundle kernelBundleInput =
sycl::get_kernel_bundle<sycl::bundle_state::input>(ctx, {kid});
// CHECK: urProgramCreateWithIL(
// CHECK: urProgramBuildExp(
// CHECK: urProgramBuild(
auto KernelBundleExe1 = build(kernelBundleInput, {dev1, dev2});
// CHECK: urProgramCreateWithIL(
// CHECK: urProgramBuildExp(
// CHECK: urProgramBuild(
auto KernelBundleExe2 = build(kernelBundleInput, {dev2, dev3});
// No other program creation calls are expected.
// CHECK-NOT: urProgramCreateWithIL(
Expand Down
18 changes: 9 additions & 9 deletions sycl/test-e2e/ProgramManager/multi_device_bundle/device_libs_and_caching.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,18 +73,18 @@ int main() {
// Create and compile the program for required device libraries (2 of them
// in this case).
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCreateWithIL(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompileExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompile(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCreateWithIL(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompileExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompile(

// Compile the main program
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompileExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompile(

// Link main program and device libraries.
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramLinkExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramLink(

// CHECK-AOT-TRACE: urProgramCreateWithBinary(
// CHECK-AOT-TRACE: urProgramBuildExp(
// CHECK-AOT-TRACE: urProgramBuild(
sycl::kernel_bundle kernelBundleExecutable =
sycl::get_kernel_bundle<sycl::bundle_state::executable>(
ctx, {dev1, dev2, dev3}, {kid});
Expand Down Expand Up @@ -131,16 +131,16 @@ int main() {
// creation is expected for device libraries as program handle already
// exists in the per-context cache.
// CHECK-SPIRV-JIT-LINK-TRACE-NOT: urProgramCreateWithIL(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompileExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompile(

// Main program will be compiled for new set of devices.
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompileExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramCompile(

// Main program will be linked with device libraries.
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramLinkExp(
// CHECK-SPIRV-JIT-LINK-TRACE: urProgramLink(

// CHECK-AOT-TRACE: urProgramCreateWithBinary(
// CHECK-AOT-TRACE: urProgramBuildExp(
// CHECK-AOT-TRACE: urProgramBuild(
sycl::kernel_bundle kernelBundleExecutableNewSet =
sycl::get_kernel_bundle<sycl::bundle_state::executable>(
ctx, {dev2, dev3, dev4}, {kid});
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,7 @@ int main() {
}

// --- Check that only a single program is built:
// CHECK: <--- urProgramBuildExp
// CHECK-NOT: <--- urProgramBuildExp
// CHECK: <--- urProgramBuild
// CHECK-NOT: <--- urProgramBuild
// --- Check that the test completed with expected results:
// CHECK: passed
8 changes: 4 additions & 4 deletions sycl/unittests/helpers/RuntimeLinkingCommon.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<ur_program_link_exp_params_t *>(pParams);
static ur_result_t redefined_urProgramLink(void *pParams) {
auto Params = *static_cast<ur_program_link_params_t *>(pParams);
unsigned ResProgram = 1;
auto Programs = *Params.pphPrograms;
for (uint32_t I = 0; I < *Params.pcount; ++I) {
Expand Down Expand Up @@ -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);
}
Loading