Skip to content
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.

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

[SYCL][Graph] Update design doc for copy queue #362

Open
wants to merge 37 commits into
base: sycl
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
37 commits
Select commit Hold shift + click to select a range
934b46f
[Driver][SYCL][NewOffload] Fix duplication of device targets (#14143)
mdtoguchi Jun 11, 2024
6ecce4f
[New offload driver][Device lib] Add SYCL device library files for al…
asudarsa Jun 11, 2024
c1b17e0
[SYCL] Enable CET for wqlibsycl-devicelib-host.a (#14135)
jinge90 Jun 12, 2024
c168f21
[UR] Fix size confusion for several device property queries (#12488)
al42and Jun 12, 2024
bdeb0ef
[SYCL][COMPAT] Added non-const image2d_max and image3d_max getters (#…
Alcpz Jun 12, 2024
d7bc4fc
[SYCL][Graph] Update L0 aspect test (#14093)
EwanC Jun 12, 2024
a497788
[SYCL][E2E] Fix CUDA include and lib paths. (#14118)
mmoadeli Jun 12, 2024
7c530e1
[UR] Bump main tag to 78d02039 (#12269)
aarongreig Jun 12, 2024
da735fe
[SYCL][COMPAT] Add math extend_v*4 to SYCLCompat (#14078)
OuadiElfarouki Jun 12, 2024
1460126
[SYCL] Remove unneeded parameter from `getOrInsertMemObjRecord` (#13807)
sergey-semenov Jun 12, 2024
bd33aaf
[E2E] Modify commands to address running on Windows. (#13682)
mmoadeli Jun 12, 2024
56f6c24
[SYCL][E2E] Fix deprecated warnings in `InorderQueue` e2e tests (#14120)
ayylol Jun 12, 2024
1a885ec
[UR] Update UR tag to include L0 loader related changes (#14109)
againull Jun 12, 2024
ae79b95
[UR] Bump main tag to b13c5e1f (#14042)
hdelan Jun 12, 2024
87f47b4
[SYCL] Remove redundant code from L0 plugin's cmake file (#14108)
againull Jun 12, 2024
3910d0c
[SYCL] Add support for key/value sorting APIs (#13942)
againull Jun 12, 2024
fe8c284
[SYCL][NewOffload][E2E] add a single test for --offload-new-driver (#…
jasonlizhengjian Jun 12, 2024
13a7b3a
[SYCL] [libdevice] Add vector overloads of ConvertBFloat16ToFINTEL an…
uditagarwal97 Jun 12, 2024
e7defab
[SYCL] Use `std::array` as storage for `sycl::vec` on device (#14130)
uditagarwal97 Jun 12, 2024
9942378
[SYCL] Adding support for missing math ops (#14132)
MaryaSharf Jun 13, 2024
e34b7ff
[Doc] Document Unified Runtime update process (#14097)
kbenzie Jun 13, 2024
f2cd2a8
[SYCL] Disable in-order queue barrier optimization while profiling (#…
sergey-semenov Jun 13, 2024
da3b5df
[SYCL] Add atomic64 aspect decoration to atomic_ref<T *> (#14052)
maksimsab Jun 13, 2024
c342a78
[SYCL] Clear cache in case of PI_ERROR_OUT_OF_HOST_MEMORY (#14119)
KornevNikita Jun 13, 2024
a5a36f8
[CI] Turn on sycl-cts/test_accessor in Nightly (#14159)
KornevNikita Jun 13, 2024
957f762
[GHA] Uplift Linux IGC Dev RT version to igc-dev-480f8b6 (#14155)
bb-sycl Jun 13, 2024
8eff95c
[SYCL] Fix FloatVecToBF16Vec build (#14161)
npmiller Jun 13, 2024
32911b2
Bump braces from 3.0.2 to 3.0.3 in /mlir/utils/vscode (#14144)
dependabot[bot] Jun 13, 2024
4acca90
[CLC][AMDGPU] Refactor fence helper to process order semantic explici…
GeorgeWeb Jun 13, 2024
c2e5529
[SYCL] Re-enable `Basic/barrier_order.cpp` (#14154)
aelovikov-intel Jun 13, 2024
4e41992
[SYCL][ESIMD][E2E] Fix rotate.cpp on Windows (#14152)
sarnex Jun 13, 2024
f9fd95e
[Driver][SYCL][NewOffloadModel] Incorporate -device settings for GPU …
mdtoguchi Jun 13, 2024
73cf85d
[SYCL][COMPAT] Add math extend_vcompare[2/4] to SYCLCompat (#14079)
OuadiElfarouki Jun 14, 2024
579484f
[UR][L0] Maintain Lock of Queue while syncing the Last Command Event …
nrspruit Jun 14, 2024
19052da
[SYCL][E2E] Use callable device selector in `FilterSelector` e2e test…
ayylol Jun 14, 2024
090c9aa
[SYCL][Graph] Update design doc for copy optimization and add test
mfrancepillois Mar 15, 2024
01b1582
Update sycl/plugins/unified_runtime/CMakeLists.txt
EwanC Jun 14, 2024
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
50 changes: 40 additions & 10 deletions clang/lib/Driver/Driver.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1182,6 +1182,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
// of -fsycl*target options passed
Arg *SYCLTargetsValues = SYCLTargets;
if (SYCLTargetsValues) {
llvm::StringSet<> SYCLTriples;
if (SYCLTargetsValues->getNumValues()) {

// Multiple targets are currently not supported when using
Expand Down Expand Up @@ -1220,15 +1221,40 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
const ToolChain *HostTC =
C.getSingleOffloadToolChain<Action::OFK_Host>();
llvm::Triple HostTriple = HostTC->getTriple();
UniqueSYCLTriplesVec.push_back(HostTriple);
SYCLTriples.insert(HostTriple.normalize());
continue;
}

if (!isValidSYCLTriple(MakeSYCLDeviceTriple(UserTargetName))) {
llvm::Triple DeviceTriple(MakeSYCLDeviceTriple(UserTargetName));
if (!isValidSYCLTriple(DeviceTriple)) {
Diag(clang::diag::err_drv_invalid_sycl_target) << Val;
continue;
}

// For any -fsycl-targets=spir64_gen additions, we will scan the
// additional -X* options for potential -device settings. These
// need to be added as a known Arch to the packager.
if (DeviceTriple.isSPIRAOT() && Arch.empty() &&
DeviceTriple.getSubArch() == llvm::Triple::SPIRSubArch_gen) {
const ToolChain *HostTC =
C.getSingleOffloadToolChain<Action::OFK_Host>();
auto DeviceTC = std::make_unique<toolchains::SYCLToolChain>(
*this, DeviceTriple, *HostTC, C.getInputArgs());
assert(DeviceTC && "Device toolchain not defined.");
ArgStringList TargetArgs;
DeviceTC->TranslateBackendTargetArgs(DeviceTC->getTriple(),
C.getInputArgs(), TargetArgs);
// Look for -device <string> and use that as the known arch to
// be associated with the current spir64_gen entry. Grab the
// right most entry.
for (int i = TargetArgs.size() - 2; i >= 0; --i) {
if (StringRef(TargetArgs[i]) == "-device") {
Arch = TargetArgs[i + 1];
break;
}
}
}

// Make sure we don't have a duplicate triple.
std::string NormalizedName = MakeSYCLDeviceTriple(Val).normalize();
auto Duplicate = FoundNormalizedTriples.find(NormalizedName);
Expand All @@ -1241,11 +1267,16 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
// Store the current triple so that we can check for duplicates in
// the following iterations.
FoundNormalizedTriples[NormalizedName] = Val;
llvm::Triple DeviceTriple(MakeSYCLDeviceTriple(UserTargetName));
UniqueSYCLTriplesVec.push_back(DeviceTriple);
SYCLTriples.insert(DeviceTriple.normalize());
if (!Arch.empty())
DerivedArchs[DeviceTriple.getTriple()].insert(Arch);
}
if (!SYCLTriples.empty()) {
for (const auto &SYCLTriple : SYCLTriples) {
llvm::Triple Triple(SYCLTriple.getKey());
UniqueSYCLTriplesVec.push_back(Triple);
}
}
addSYCLDefaultTriple(C, UniqueSYCLTriplesVec);
} else
Diag(clang::diag::warn_drv_empty_joined_argument)
Expand Down Expand Up @@ -5031,17 +5062,16 @@ class OffloadingActionBuilder final {
}

// By default, we produce an action for each device arch.
auto TC = ToolChains.begin();
for (Action *&A : SYCLDeviceActions) {
if ((*TC)->getTriple().isNVPTX() && CurPhase >= phases::Backend) {
for (auto TargetActionInfo :
llvm::zip(SYCLDeviceActions, SYCLTargetInfoList)) {
auto &TargetInfo = std::get<1>(TargetActionInfo);
if (TargetInfo.TC->getTriple().isNVPTX() && CurPhase >= phases::Backend)
// For CUDA, stop to emit LLVM IR so it can be linked later on.
++TC;
continue;
}

Action *&A = std::get<0>(TargetActionInfo);
A = C.getDriver().ConstructPhaseAction(C, Args, CurPhase, A,
AssociatedOffloadKind);
++TC;
}

return ABRT_Success;
Expand Down
63 changes: 22 additions & 41 deletions clang/lib/Driver/ToolChains/Clang.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -11005,13 +11005,6 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
if (Args.hasArg(options::OPT_v))
CmdArgs.push_back("--wrapper-verbose");

// TODO(NOM2): Pass following options to clang-linker-wrapper.
// Please refer to sycl/doc/design/OffloadDesign.md for details.
// sycl-device-libraries
// sycl-device-library-location
// sycl-post-link-options
// llvm-spirv-options

if (const Arg *A = Args.getLastArg(options::OPT_g_Group)) {
if (!A->getOption().matches(options::OPT_g0))
CmdArgs.push_back("--device-debug");
Expand Down Expand Up @@ -11044,12 +11037,14 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
// Add any SYCL offloading specific options to the clang-linker-wrapper
if (C.hasOffloadToolChain<Action::OFK_SYCL>()) {
// -sycl-device-libraries=<comma separated list> contains all of the SYCL
// device specific libraries that are needed. This provides the list of
// files file only.
// TODO: This generic list will be populated with only device binaries
// for spir/spirv. Other targets (AOT and others) can represent a different
// set of device libraries. We will cross that bridge when we begin to
// enable the other possible targets.
// device specific libraries that are needed. This generic list will be
// populated with device binaries for all target triples in the current
// compilation flow.

// Create a comma separated list to pass along to the linker wrapper.
SmallString<256> LibList;
// TODO: TargetTriple should not be used here for creating linker wrapper
// options. It should also not be passed to the linker wrapper.
llvm::Triple TargetTriple;
auto ToolChainRange = C.getOffloadToolChains<Action::OFK_SYCL>();
for (auto &I :
Expand All @@ -11058,38 +11053,24 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
if (TC->getTriple().isSPIROrSPIRV() &&
TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) {
TargetTriple = TC->getTriple();
break;
SmallVector<std::string, 8> SYCLDeviceLibs;
bool IsSPIR = TargetTriple.isSPIROrSPIRV();
bool IsSpirvAOT = TargetTriple.isSPIRAOT();
bool UseJitLink =
IsSPIR &&
Args.hasFlag(options::OPT_fsycl_device_lib_jit_link,
options::OPT_fno_sycl_device_lib_jit_link, false);
bool UseAOTLink = IsSPIR && (IsSpirvAOT || !UseJitLink);
SYCLDeviceLibs = SYCL::getDeviceLibraries(C, TargetTriple, UseAOTLink);
for (const auto &AddLib : SYCLDeviceLibs) {
if (LibList.size() > 0)
LibList += ",";
LibList += AddLib;
}
}
}
// Pass the device triple to the linker wrapper tool for SYCL offload.
// Only spir64 or spirv64 is currently passed.
// TODO(NOM1): Support target triples in a more generic way.
// TODO(NOM3): Investigate why passing spirv64-unknown-unknown does not
// work.
if (TargetTriple.isSPIR())
CmdArgs.push_back("--triple=spir64");
else if (TargetTriple.isSPIRV())
CmdArgs.push_back("--triple=spirv64");

SmallVector<std::string, 8> SYCLDeviceLibs;
auto IsSPIR = TargetTriple.isSPIROrSPIRV();
bool IsSpirvAOT = TargetTriple.isSPIRAOT();
bool UseJitLink =
IsSPIR &&
Args.hasFlag(options::OPT_fsycl_device_lib_jit_link,
options::OPT_fno_sycl_device_lib_jit_link, false);
bool UseAOTLink = IsSPIR && (IsSpirvAOT || !UseJitLink);
SYCLDeviceLibs = SYCL::getDeviceLibraries(C, TargetTriple, UseAOTLink);
// Create a comma separated list to pass along to the linker wrapper.
SmallString<256> LibList;
for (const auto &AddLib : SYCLDeviceLibs) {
if (LibList.size() > 0)
LibList += ",";
LibList += AddLib;
}
// -sycl-device-libraries=<libs> provides a comma separate list of
// libraries to add to the device linking step.
// SYCL device libraries can be found.
if (LibList.size())
CmdArgs.push_back(
Args.MakeArgString(Twine("-sycl-device-libraries=") + LibList));
Expand Down
2 changes: 1 addition & 1 deletion clang/test/Driver/linker-wrapper-sycl-win.cpp
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
// REQUIRES: system-windows

/// Check for list of commands for standalone clang-linker-wrapper run for sycl
// RUN: clang-linker-wrapper -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.new.obj,libsycl-complex.new.obj -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-pc-windows-msvc" "--triple=spir64" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %S/Inputs/test-sycl.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s
// RUN: clang-linker-wrapper -sycl-device-library-location=%S/Inputs -sycl-device-libraries=libsycl-crt.new.obj,libsycl-complex.new.obj -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-pc-windows-msvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %S/Inputs/test-sycl.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-CMDS %s
// CHK-CMDS: "{{.*}}spirv-to-ir-wrapper.exe" {{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts=--spirv-preserve-auxdata --llvm-spirv-opts=--spirv-target-env=SPV-IR --llvm-spirv-opts=--spirv-builtin-format=global
// CHK-CMDS-NEXT: "{{.*}}llvm-link.exe" [[FIRSTLLVMLINKIN:.*]].bc -o [[FIRSTLLVMLINKOUT:.*]].bc --suppress-warnings
// CHK-CMDS-NEXT: "{{.*}}llvm-link.exe" -only-needed [[FIRSTLLVMLINKOUT]].bc {{.*}}.bc {{.*}}.bc -o [[SECONDLLVMLINKOUT:.*]].bc --suppress-warnings
Expand Down
4 changes: 2 additions & 2 deletions clang/test/Driver/sycl-fno-libspirv-warn.cpp
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
/// Test that appropriate warnings are output when -fno-sycl-libspirv is used.

// RUN: not %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda,amdgcn-amd-amdhsa -fno-sycl-libspirv %s -### 2>&1 | FileCheck %s
// CHECK: warning: '-fno-sycl-libspirv' should not be used with target 'nvptx64-nvidia-cuda'; libspirv is required for correct behavior [-Wno-libspirv-hip-cuda]
// CHECK: warning: '-fno-sycl-libspirv' should not be used with target 'amdgcn-amd-amdhsa'; libspirv is required for correct behavior [-Wno-libspirv-hip-cuda]
// CHECK-DAG: warning: '-fno-sycl-libspirv' should not be used with target 'nvptx64-nvidia-cuda'; libspirv is required for correct behavior [-Wno-libspirv-hip-cuda]
// CHECK-DAG: warning: '-fno-sycl-libspirv' should not be used with target 'amdgcn-amd-amdhsa'; libspirv is required for correct behavior [-Wno-libspirv-hip-cuda]
// RUN: %clangxx -fsycl -fsycl-targets=spir64-unknown-unknown -fno-sycl-libspirv %s -### 2>&1 | FileCheck --check-prefix=CHECK-SPIR64 %s
// CHECK-SPIR64: ignoring '-fno-sycl-libspirv' option as it is not currently supported for target 'spir64-unknown-unknown' [-Woption-ignored]
2 changes: 1 addition & 1 deletion clang/test/Driver/sycl-linker-wrapper-image.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
// RUN: %clang -cc1 -fsycl-is-device -disable-llvm-passes -triple=spir64-unknown-unknown %s -emit-llvm-bc -o %t.device.bc
// RUN: clang-offload-packager -o %t.fat --image=file=%t.device.bc,kind=sycl,triple=spir64-unknown-unknown
// RUN: %clang -cc1 %s -triple=x86_64-unknown-linux-gnu -emit-obj -o %t.o -fembed-offload-object=%t.fat
// RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu --triple=spir64 \
// RUN: clang-linker-wrapper --print-wrapped-module --host-triple=x86_64-unknown-linux-gnu \
// RUN: -sycl-device-library-location=%S/Inputs -sycl-post-link-options="-split=auto -symbols" \
// RUN: %t.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s

Expand Down
Loading
Loading