Skip to content

[SYCL-MLIR] pulldown (20230531) #9666

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

Merged
merged 24 commits into from
May 31, 2023
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
24 commits
Select commit Hold shift + click to select a range
b682bc9
Adjust CODEOWNERS for OffloadBundler.{h,cpp} files (#9618)
maksimsab May 28, 2023
37aa84b
[SYCL] Improve the accuracy of host `sycl::cospi` (#9575)
0x12CC May 30, 2023
a410510
[SYCL] Remove undefined behavior in host-side abs_diff (#9627)
steffenlarsen May 30, 2023
be080bf
[SYCL][E2E] Add `sycl::free` to root group tests (#9639)
0x12CC May 30, 2023
4ed3676
[Driver][SYCL] Improve FPGA archive device unbundling with AOCO (#9572)
maksimsab May 30, 2023
c40baa6
[SYCL][ESIMD][Emulator] Implement device query for 64 bit atomics (#9…
turinevgeny May 30, 2023
9562eb3
[SYCL][Matrix] Fix TK size for tf32 test case. (#9531)
arnamoy10 May 30, 2023
4bf5423
[SYCL][CUDA][HIP] Add support for sycl::aspect::ext_intel_pci_address…
al42and May 30, 2023
4e9ab7a
[SYCL][Test E2E] Remove sycl_be/target_devices params support (#9606)
aelovikov-intel May 30, 2023
034ac09
[SYCL][Doc] Add draft of sycl_ext_oneapi_prefetch (#9282)
Pennycook May 30, 2023
f40e48d
[SYCL][UR] Update UR loader (#9637)
May 30, 2023
98cdf4c
[SYCL][PI] Fix `PI_KERNEL_MAX_SUB_GROUP_SIZE` query for devices witho…
0x12CC May 30, 2023
eafd768
Add support for fpbuiltin accuracy lookup (#9167)
May 30, 2023
a3a1eee
[SYCL][ABI Breaking] CG/handler refactoring to reduce boilerplate (#9…
aelovikov-intel May 30, 2023
ffdcd36
[SYCL] Add tests that submit kernels from shared library (#9626)
aelovikov-intel May 30, 2023
8cceeb5
[OffloadBundler] Fix build with clang-10 (#9649)
aelovikov-intel May 30, 2023
566f9ce
[SYCL][UR][L0] Fix casting in UR for boolean (#9654)
May 31, 2023
1251517
[SYCL] Print supported SG sizes in sycl-ls --verbose (#9481)
aelovikov-intel May 31, 2023
399442d
[Driver] Fix FPGAEmulationMode naming confusion (#9537)
srividya-sundaram May 31, 2023
fd45980
[SYCL] Modifies accessor::ConcreteASPtrType to be `const` for `readon…
mmoadeli May 31, 2023
26ac54b
[SYCL][Matrix] Allow get_coord() tests to pass on GPU and fail on CPU…
arnamoy10 May 31, 2023
40ac9cc
[SYCL][Test] Remove check_has.cpp dependency on metadata order (#9642)
jinge90 May 31, 2023
7358559
Merge remote-tracking branch 'intel_llvm/sycl' into sycl_mlir_pulldown
sys-ce-bb May 31, 2023
3597fbc
Revert "[SYCL] Modifies accessor::ConcreteASPtrType to be `const` for…
whitneywhtsang May 31, 2023
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
30 changes: 19 additions & 11 deletions clang/include/clang/Driver/Driver.h
Original file line number Diff line number Diff line change
Expand Up @@ -718,6 +718,25 @@ class Driver {
return IsOffload ? OffloadLTOMode : LTOMode;
}

// FPGA Offload Modes.
enum DeviceMode {
UnsetDeviceMode,
FPGAHWMode,
FPGAEmulationMode
} OffloadCompileMode = UnsetDeviceMode;

bool IsFPGAHWMode() const { return OffloadCompileMode == FPGAHWMode; }

bool IsFPGAEmulationMode() const {
return OffloadCompileMode == FPGAEmulationMode;
}

void setOffloadCompileMode(DeviceMode ModeValue) {
OffloadCompileMode = ModeValue;
}

DeviceMode getOffloadCompileMode() { return OffloadCompileMode; }

private:

/// Tries to load options from configuration files.
Expand Down Expand Up @@ -793,13 +812,6 @@ class Driver {
bool UseNewOffloadingDriver = false;
void setUseNewOffloadingDriver() { UseNewOffloadingDriver = true; }

/// FPGA Emulation Mode. By default, this is true due to the fact that
/// an external option setting is required to target hardware.
bool FPGAEmulationMode = true;
void setFPGAEmulationMode(bool IsEmulation) {
FPGAEmulationMode = IsEmulation;
}

/// The inclusion of the default SYCL device triple is dependent on either
/// the discovery of an existing object/archive that contains the device code
/// or if a user explicitly turns this on with -fsycl-add-spirv.
Expand Down Expand Up @@ -885,10 +897,6 @@ class Driver {
return FPGATempDepFiles[FileName];
}

/// isFPGAEmulationMode - Compilation mode is determined to be used for
/// FPGA Emulation. This is only used for SYCL offloading to FPGA device.
bool isFPGAEmulationMode() const { return FPGAEmulationMode; };

/// isSYCLDefaultTripleImplied - The default SYCL triple (spir64) has been
/// added or should be added given proper criteria.
bool isSYCLDefaultTripleImplied() const { return SYCLDefaultTripleImplied; };
Expand Down
3 changes: 3 additions & 0 deletions clang/include/clang/Driver/OffloadBundler.h
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,9 @@ class OffloadBundlerConfig {
std::vector<std::string> TargetNames;
std::vector<std::string> InputFileNames;
std::vector<std::string> OutputFileNames;

// List of excluded target names from unbundling.
std::vector<std::string> ExcludedTargetNames;
};

class OffloadBundler {
Expand Down
20 changes: 9 additions & 11 deletions clang/lib/Driver/Driver.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1841,9 +1841,12 @@ Compilation *Driver::BuildCompilation(ArrayRef<const char *> ArgList) {
static_cast<const toolchains::SYCLToolChain *>(TI->second);
FPGATC->TranslateBackendTargetArgs(FPGATC->getTriple(), *TranslatedArgs,
TargetArgs);
// By default, FPGAEmulationMode is true due to the fact that
// an external option setting is required to target hardware.
setOffloadCompileMode(FPGAEmulationMode);
for (StringRef ArgString : TargetArgs) {
if (ArgString.equals("-hardware") || ArgString.equals("-simulation")) {
setFPGAEmulationMode(false);
setOffloadCompileMode(FPGAHWMode);
break;
}
}
Expand Down Expand Up @@ -6275,7 +6278,7 @@ class OffloadingActionBuilder final {
FPGAOutType = (A->getValue() == StringRef("early"))
? types::TY_FPGA_AOCR
: types::TY_FPGA_AOCX;
if (C.getDriver().isFPGAEmulationMode())
if (C.getDriver().IsFPGAEmulationMode())
FPGAOutType = (A->getValue() == StringRef("early"))
? types::TY_FPGA_AOCR_EMU
: types::TY_FPGA_AOCX;
Expand Down Expand Up @@ -6492,7 +6495,6 @@ class OffloadingActionBuilder final {
DerivedArgList &Args) {
std::string InputName = InputArg->getAsString(Args);
const Driver &D = C.getDriver();
bool IsFPGAEmulation = D.isFPGAEmulationMode();
// Only check for FPGA device information when using fpga SubArch.
if (A->getType() == types::TY_Object && isObjectFile(InputName))
return true;
Expand Down Expand Up @@ -6520,12 +6522,13 @@ class OffloadingActionBuilder final {
{types::TY_FPGA_AOCR_EMU, true}};
for (const auto &ArchiveType : FPGAAOCTypes) {
bool BinaryFound = hasFPGABinary(C, InputName, ArchiveType.first);
if (BinaryFound && ArchiveType.second == IsFPGAEmulation) {
if (BinaryFound && ArchiveType.second == D.IsFPGAEmulationMode()) {
// Binary matches check and emulation type, we keep this one.
A = C.MakeAction<InputAction>(*InputArg, ArchiveType.first);
return true;
}
ArchiveTypeMismatch(BinaryFound && ArchiveType.second != IsFPGAEmulation);
ArchiveTypeMismatch(BinaryFound &&
ArchiveType.second == D.IsFPGAHWMode());
}
return true;
}
Expand Down Expand Up @@ -6871,13 +6874,8 @@ class OffloadingActionBuilder final {
// unbundling for FPGA AOT static lib usage. Uses FPGA aoco type to
// differentiate if aoco unbundling is needed. Unbundling of aoco is
// not needed for emulation, as these are treated as regular archives.
if (!C.getDriver().isFPGAEmulationMode())
if (C.getDriver().IsFPGAHWMode())
unbundleStaticLib(types::TY_FPGA_AOCO, LA);
// Do not unbundle any AOCO archive as a regular archive when we are
// in FPGA Hardware/Simulation mode.
if (!C.getDriver().isFPGAEmulationMode() &&
hasFPGABinary(C, LA.str(), types::TY_FPGA_AOCO))
continue;
unbundleStaticLib(types::TY_Archive, LA);
}
}
Expand Down
90 changes: 90 additions & 0 deletions clang/lib/Driver/OffloadBundler.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,7 @@
#include <set>
#include <string>
#include <system_error>
#include <unordered_set>
#include <utility>

using namespace llvm;
Expand All @@ -71,6 +72,8 @@ using namespace clang;
/// Section name which holds target symbol names.
#define SYMBOLS_SECTION_NAME ".tgtsym"

#define DEBUG_TYPE "clang-offload-bundler"

OffloadTargetInfo::OffloadTargetInfo(const StringRef Target,
const OffloadBundlerConfig &BC)
: BundlerConfig(BC) {
Expand Down Expand Up @@ -1015,6 +1018,10 @@ class ArchiveFileHandler final : public FileHandler {
.Case("a", OutputType::Archive)
.Default(OutputType::Unknown);

// Set contains indexes of Children that should be skipped during
// unbundling.
std::unordered_set<size_t> ExcludedChildIndexes;

public:
ArchiveFileHandler(const OffloadBundlerConfig &BC) : BundlerConfig(BC) {}
~ArchiveFileHandler() = default;
Expand All @@ -1029,8 +1036,10 @@ class ArchiveFileHandler final : public FileHandler {
Ar = std::move(*ArOrErr);

// Read all children.
ssize_t ChildIndex = -1;
Error Err = Error::success();
for (auto &C : Ar->children(Err)) {
++ChildIndex;
auto BinOrErr = C.getAsBinary();
if (!BinOrErr) {
if (auto Err = isNotObjectErrorInvalidFileType(BinOrErr.takeError()))
Expand All @@ -1042,6 +1051,16 @@ class ArchiveFileHandler final : public FileHandler {
if (!Bin->isObject())
continue;

auto CheckOrErr = CheckIfObjectFileContainsExcludedTargets(C);
if (!CheckOrErr)
return CheckOrErr.takeError();

if (*CheckOrErr) {
LLVM_DEBUG(outs() << "Add child to ban list. Index: " << ChildIndex
<< "\n");
ExcludedChildIndexes.emplace(ChildIndex);
}

auto Obj = std::unique_ptr<ObjectFile>(cast<ObjectFile>(Bin.release()));
auto Buf = MemoryBuffer::getMemBuffer(Obj->getMemoryBufferRef(), false);

Expand Down Expand Up @@ -1099,7 +1118,14 @@ class ArchiveFileHandler final : public FileHandler {

// Read all children.
Error Err = Error::success();
ssize_t ChildIndex = -1;
for (auto &C : Ar->children(Err)) {
++ChildIndex;
if (ExcludedChildIndexes.count(ChildIndex)) {
LLVM_DEBUG(outs() << "Skip Child. Index: " << ChildIndex << "\n");
continue;
}

auto BinOrErr = C.getAsBinary();
if (!BinOrErr) {
if (auto Err = isNotObjectErrorInvalidFileType(BinOrErr.takeError()))
Expand Down Expand Up @@ -1215,6 +1241,70 @@ class ArchiveFileHandler final : public FileHandler {
Error WriteBundle(raw_fd_ostream &OS, MemoryBuffer &Input) override {
llvm_unreachable("unsupported for the ArchiveFileHandler");
}

private:
// NOTE: mostly a copy-paste of ReadHeader method.
Expected<std::vector<std::string>>
ReadTargetsFromChild(const Archive::Child &C) {
Expected<std::unique_ptr<Binary>> BinOrErr = C.getAsBinary();
if (!BinOrErr)
return BinOrErr.takeError();

std::unique_ptr<Binary> &Bin = BinOrErr.get();
auto Obj = std::unique_ptr<ObjectFile>(cast<ObjectFile>(Bin.release()));
std::unique_ptr<MemoryBuffer> Buf =
MemoryBuffer::getMemBuffer(Obj->getMemoryBufferRef(), false);
ObjectFileHandler OFH(std::move(Obj), BundlerConfig);
if (Error Err = OFH.ReadHeader(*Buf))
return {std::move(Err)};
Expected<std::optional<StringRef>> NameOrErr = OFH.ReadBundleStart(*Buf);
if (!NameOrErr)
return NameOrErr.takeError();

std::vector<std::string> Targets;
while (*NameOrErr) {
if (*NameOrErr)
Targets.emplace_back((**NameOrErr).str());
NameOrErr = OFH.ReadBundleStart(*Buf);
if (!NameOrErr)
return NameOrErr.takeError();
}

return Targets;
}

bool CheckIfTargetIsExcluded(StringRef Triple) {
// NOTE: "-sycldevice" Triple component has been deprecated.
// However, it still can be met in libraries that have been compiled before
// deprecation. For example, here Triple might be the following:
// sycl-fpga_aoco-intel-unknown-sycldevice
//
// The workaround is to strip this Triple component if it is present.
Triple.consume_back("-sycldevice");
const auto &ExcludedTargetNames = BundlerConfig.ExcludedTargetNames;
auto It = std::find(ExcludedTargetNames.begin(), ExcludedTargetNames.end(),
Triple);
return It != ExcludedTargetNames.end();
}

// Function reads targets from Child and checks whether one of Targets
// is in Excluded list.
Expected<bool>
CheckIfObjectFileContainsExcludedTargets(const Archive::Child &C) {
if (BundlerConfig.ExcludedTargetNames.empty())
return false;

auto TargetNamesOrErr = ReadTargetsFromChild(C);
if (!TargetNamesOrErr)
return TargetNamesOrErr.takeError();

auto TargetNames = TargetNamesOrErr.get();
for (const auto &TargetName : TargetNames)
if (CheckIfTargetIsExcluded(TargetName))
return true;

return false;
}
};

/// Return an appropriate object file handler. We use the specific object
Expand Down
30 changes: 26 additions & 4 deletions clang/lib/Driver/ToolChains/Clang.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -9114,15 +9114,37 @@ void OffloadBundler::ConstructJobMultipleOutputs(
TypeArg = "o";

bool HasSPIRTarget = false;
bool HasFPGATarget = false;
auto SYCLTCRange = C.getOffloadToolChains<Action::OFK_SYCL>();
for (auto TI = SYCLTCRange.first, TE = SYCLTCRange.second; TI != TE; ++TI)
HasSPIRTarget |= TI->second->getTriple().isSPIR();
for (auto TI = SYCLTCRange.first, TE = SYCLTCRange.second; TI != TE; ++TI) {
llvm::Triple TT(TI->second->getTriple());
if (TT.isSPIR()) {
HasSPIRTarget = true;
if (TT.getSubArch() == llvm::Triple::SPIRSubArch_fpga)
HasFPGATarget = true;
}
}
if (InputType == types::TY_Archive && HasSPIRTarget)
TypeArg = "aoo";

// Get the type.
CmdArgs.push_back(TCArgs.MakeArgString(Twine("-type=") + TypeArg));

// For FPGA Archives that contain AOCO in them, we only want to unbundle
// the objects from the archive that do not have AOCO associated in that
// specific object. Only do this when in hardware mode.
if (InputType == types::TY_Archive && HasFPGATarget && !IsFPGADepUnbundle &&
!IsFPGADepLibUnbundle && C.getDriver().IsFPGAHWMode()) {
llvm::Triple TT;
TT.setArchName(types::getTypeName(types::TY_FPGA_AOCO));
TT.setVendorName("intel");
TT.setOS(getToolChain().getTriple().getOS());
SmallString<128> ExcludedTargets("-excluded-targets=");
ExcludedTargets += "sycl-";
ExcludedTargets += TT.normalize();
CmdArgs.push_back(TCArgs.MakeArgString(ExcludedTargets));
}

// Get the targets.
SmallString<128> Triples;
Triples += "-targets=";
Expand Down Expand Up @@ -9289,7 +9311,7 @@ void OffloadWrapper::ConstructJob(Compilation &C, const JobAction &JA,
auto *A = C.getInputArgs().getLastArg(options::OPT_fsycl_link_EQ);
bool Early = (A->getValue() == StringRef("early"));
FPGAArch += Early ? "aocr" : "aocx";
if (C.getDriver().isFPGAEmulationMode() && Early)
if (C.getDriver().IsFPGAEmulationMode() && Early)
FPGAArch += "_emu";
TT.setArchName(FPGAArch);
TT.setVendorName("intel");
Expand Down Expand Up @@ -9658,7 +9680,7 @@ void SPIRVTranslator::ConstructJob(Compilation &C, const JobAction &JA,
",+SPV_INTEL_fpga_argument_interfaces"
",+SPV_INTEL_fpga_invocation_pipelining_attributes";
ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
if (!C.getDriver().isFPGAEmulationMode())
if (C.getDriver().IsFPGAHWMode())
// Enable several extensions on FPGA H/W exclusively
ExtArg += ",+SPV_INTEL_usm_storage_classes,+SPV_INTEL_runtime_aligned"
",+SPV_INTEL_fpga_cluster_attributes,+SPV_INTEL_loop_fuse"
Expand Down
4 changes: 2 additions & 2 deletions clang/lib/Driver/ToolChains/SYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -445,7 +445,7 @@ void SYCL::fpga::BackendCompiler::ConstructJob(

// When performing emulation compilations for FPGA AOT, we want to use
// opencl-aot instead of aoc.
if (C.getDriver().isFPGAEmulationMode()) {
if (C.getDriver().IsFPGAEmulationMode()) {
constructOpenCLAOTCommand(C, JA, Output, Inputs, Args);
return;
}
Expand Down Expand Up @@ -850,7 +850,7 @@ SYCLToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
}
}
// Strip out -O0 for FPGA Hardware device compilation.
if (!getDriver().isFPGAEmulationMode() &&
if (getDriver().IsFPGAHWMode() &&
getTriple().getSubArch() == llvm::Triple::SPIRSubArch_fpga)
DAL->eraseArg(options::OPT_O0);

Expand Down
Binary file not shown.
36 changes: 36 additions & 0 deletions clang/test/Driver/clang-offload-bundler-exclude.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,36 @@
// This test prepares Archive input for clang-offload-bundler
// and checks -exclude-target command line option.
// Option should exclude fat_device_aoco object file.

// UNSUPPORTED: system-windows

// The test uses assembled archive file fatlib.a.
// The assembly algorithm is the following:
// echo "DUMMY IR FILE" > device
// echo "DUMMY IR2 FILE" > device2
// echo "DUMMY AOCO FILE" > aoco
// echo "DUMMY HOST FILE" > host
// echo "DUMMY HOST2 FILE" > host2
// # Wrap and compile objects
// clang-offload-wrapper -o=device.bc -host=x86_64-unknown-linux-gnu -target=spir64 -kind=sycl device
// clang-offload-wrapper -o=device2.bc -host=x86_64-unknown-linux-gnu -target=spir64 -kind=sycl device2
// clang-offload-wrapper -o=aoco.bc -host=x86_64-unknown-linux-gnu -target=spir64 -kind=sycl aoco
// clang-offload-wrapper -o=host.bc -host=x86_64-unknown-linux-gnu -target=spir64 -kind=sycl host
// clang-offload-wrapper -o=host2.bc -host=x86_64-unknown-linux-gnu -target=spir64 -kind=sycl host2
// llc -filetype=obj -o device.o device.bc
// llc -filetype=obj -o device2.o device2.bc
// llc -filetype=obj -o aoco.o aoco.bc
// llc -filetype=obj -o host.o host.bc
// llc -filetype=obj -o host2.o host2.bc
// # Bundle the objects
// clang-offload-bundler -input=device.o -input=host.o -output=fat_device.o -targets=sycl-spir64_fpga-unknown-unknown,host-x86_64-unknown-linux-gnu -type=o
// clang-offload-bundler -input=device2.o -input=aoco.o -input=host2.o -output=fat_device_aoco.o -targets=sycl-spir64_fpga-unknown-unknown,sycl-fpga_aoco-intel-unknown,host-x86_64-unknown-linux-gnu -type=o
// # Create the archive
// ar cr fatlib.a fat_device.o fat_device_aoco.o


// Unbundle archive
// RUN: clang-offload-bundler -type=aoo -excluded-targets=sycl-fpga_aoco-intel-unknown -targets=sycl-spir64_fpga-unknown-unknown -input=%S/Inputs/clang-offload-bundler-exclude/fatlib.a -output=%t-my_output.txt -unbundle -allow-missing-bundles

// Check that output of unbundling doesn't contain content of device2
// RUN: cat %t-my_output.txt | xargs cat | strings | not grep "DUMMY IR2"
Loading