Skip to content

Commit 1a64242

Browse files
AllanZyneomarahmed1111
authored andcommitted
Merge branch 'sycl' into review/yang/dsan_nullpointer
2 parents dd06c90 + 18081b9 commit 1a64242

File tree

111 files changed

+1814
-414
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

111 files changed

+1814
-414
lines changed

.github/CODEOWNERS

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -97,14 +97,14 @@ invoke_simd/ @intel/dpcpp-esimd-reviewers
9797
InvokeSimd/ @intel/dpcpp-esimd-reviewers
9898
sycl/test-e2e/InvokeSimd/ @intel/dpcpp-esimd-reviewers
9999

100-
# dev-igc driver update
101-
devops/dependencies-igc-dev.json @intel/sycl-matrix-reviewers @intel/dpcpp-esimd-reviewers @intel/dpcpp-devops-reviewers
102-
103100
# DevOps configs
104101
.github/ @intel/dpcpp-devops-reviewers
105102
buildbot/ @intel/dpcpp-devops-reviewers
106103
devops/ @intel/dpcpp-devops-reviewers
107104

105+
# dev-igc driver update
106+
devops/dependencies-igc-dev.json @intel/sycl-matrix-reviewers @intel/dpcpp-esimd-reviewers @intel/dpcpp-devops-reviewers
107+
108108
# Kernel fusion JIT compiler
109109
sycl-jit/ @intel/dpcpp-kernel-fusion-reviewers
110110
sycl/doc/design/KernelFusionJIT.md @intel/dpcpp-kernel-fusion-reviewers

.github/workflows/sycl-linux-precommit-aws.yml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -66,7 +66,7 @@ jobs:
6666
with:
6767
name: CUDA E2E
6868
runner: '["aws_cuda-${{ github.event.workflow_run.id }}-${{ github.event.workflow_run.run_attempt }}"]'
69-
image: ghcr.io/intel/llvm/ubuntu2204_build:latest-0300ac924620a51f76c4929794637b82790f12ab
69+
image: ghcr.io/intel/llvm/ubuntu2204_build:latest
7070
image_options: -u 1001 --gpus all --cap-add SYS_ADMIN --env NVIDIA_DISABLE_REQUIRE=1
7171
target_devices: ext_oneapi_cuda:gpu
7272
# No idea why but that seems to work and be in sync with the main

.github/workflows/sycl-linux-precommit.yml

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,6 @@ jobs:
4646
build_artifact_suffix: "default"
4747
build_cache_suffix: "default"
4848
changes: ${{ needs.detect_changes.outputs.filters }}
49-
build_image: "ghcr.io/intel/llvm/ubuntu2204_build:latest-0300ac924620a51f76c4929794637b82790f12ab"
5049

5150
determine_arc_tests:
5251
name: Decide which Arc tests to run

.github/workflows/sycl-post-commit.yml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -60,7 +60,7 @@ jobs:
6060
reset_intel_gpu: true
6161
- name: AMD/HIP
6262
runner: '["Linux", "amdgpu"]'
63-
image: ghcr.io/intel/llvm/ubuntu2204_build:latest-0300ac924620a51f76c4929794637b82790f12ab
63+
image: ghcr.io/intel/llvm/ubuntu2204_build:latest
6464
image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
6565
target_devices: ext_oneapi_hip:gpu
6666
reset_intel_gpu: false

clang/lib/Driver/ToolChains/SYCL.cpp

Lines changed: 9 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -592,16 +592,15 @@ const char *SYCL::Linker::constructLLVMLinkCommand(
592592
NewLibPostfix = ".new.obj";
593593
std::string FileName = this->getToolChain().getInputFilename(II);
594594
StringRef InputFilename = llvm::sys::path::filename(FileName);
595-
if (IsNVPTX || IsSYCLNativeCPU) {
596-
// Linking SYCL Device libs requires libclc as well as libdevice
597-
if ((InputFilename.find("libspirv") != InputFilename.npos ||
598-
InputFilename.find("libdevice") != InputFilename.npos))
599-
return true;
600-
if (IsNVPTX) {
601-
LibPostfix = ".cubin";
602-
NewLibPostfix = ".new.cubin";
603-
}
604-
}
595+
// NativeCPU links against libclc (libspirv)
596+
if (IsSYCLNativeCPU && InputFilename.contains("libspirv"))
597+
return true;
598+
// NVPTX links against our libclc (libspirv), our libdevice (devicelib),
599+
// and the CUDA libdevice
600+
if (IsNVPTX && (InputFilename.starts_with("devicelib-") ||
601+
InputFilename.contains("libspirv") ||
602+
InputFilename.contains("libdevice")))
603+
return true;
605604
StringRef LibSyclPrefix("libsycl-");
606605
if (!InputFilename.starts_with(LibSyclPrefix) ||
607606
!InputFilename.ends_with(LibPostfix) ||
Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
// Check that we correctly determine that the final link command links
2+
// devicelibs together, as far as the driver is concerned. This results in the
3+
// -only-needed flag.
4+
//
5+
// Note we check the names of the various device libraries because that's the
6+
// logic the driver uses.
7+
8+
// Older CUDA versions had versioned libdevice files. We don't support CUDA
9+
// this old in SYCL, but we still test the driver's ability to pick out the
10+
// correctly versioned libdevice. We use Inputs/CUDA_80 which has a full set of
11+
// libdevice files.
12+
// RUN: %clang -### -fsycl -fsycl-targets=nvptx64-nvidia-cuda \
13+
// RUN: -Xsycl-target-backend --cuda-gpu-arch=sm_30 \
14+
// RUN: --sysroot=%S/Inputs/SYCL --cuda-path=%S/Inputs/CUDA_80/usr/local/cuda %s 2>&1 \
15+
// RUN: | FileCheck %s --check-prefixes=CHECK,LIBDEVICE30
16+
// RUN: %clang -### -fsycl -fsycl-targets=nvptx64-nvidia-cuda \
17+
// RUN: -Xsycl-target-backend --cuda-gpu-arch=sm_35 \
18+
// RUN: --sysroot=%S/Inputs/SYCL --cuda-path=%S/Inputs/CUDA_80/usr/local/cuda %s 2>&1 \
19+
// RUN: | FileCheck %s --check-prefixes=CHECK,LIBDEVICE35
20+
// RUN: %clang -### -fsycl -fsycl-targets=nvptx64-nvidia-cuda \
21+
// RUN: -Xsycl-target-backend --cuda-gpu-arch=sm_50 \
22+
// RUN: --sysroot=%S/Inputs/SYCL --cuda-path=%S/Inputs/CUDA_80/usr/local/cuda %s 2>&1 \
23+
// RUN: | FileCheck %s --check-prefixes=CHECK,LIBDEVICE50
24+
25+
// CUDA-9+ uses the same libdevice for all GPU variants
26+
// RUN: %clang -### -fsycl -fsycl-targets=nvptx64-nvidia-cuda \
27+
// RUN: -Xsycl-target-backend --cuda-gpu-arch=sm_35 \
28+
// RUN: --sysroot=%S/Inputs/SYCL --cuda-path=%S/Inputs/CUDA_90/usr/local/cuda %s 2>&1 \
29+
// RUN: | FileCheck %s --check-prefixes=CHECK,LIBDEVICE10
30+
31+
// First link command: ignored
32+
// CHECK: llvm-link
33+
34+
// CHECK: llvm-link
35+
// CHECK-SAME: -only-needed
36+
// CHECK-SAME: devicelib--cuda.bc
37+
// CHECK-SAME: libspirv-nvptx64-nvidia-cuda.bc
38+
// LIBDEVICE10-SAME: libdevice.10.bc
39+
// LIBDEVICE30-SAME: libdevice.compute_30.10.bc
40+
// LIBDEVICE35-SAME: libdevice.compute_35.10.bc
41+
// LIBDEVICE50-SAME: libdevice.compute_50.10.bc

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

Lines changed: 137 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -251,8 +251,8 @@ Expected<OffloadFile> getInputBitcodeLibrary(StringRef Input) {
251251
Image.StringData["arch"] = Arch;
252252
Image.Image = std::move(*ImageOrError);
253253

254-
std::unique_ptr<MemoryBuffer> Binary =
255-
MemoryBuffer::getMemBufferCopy(OffloadBinary::write(Image));
254+
std::unique_ptr<MemoryBuffer> Binary = MemoryBuffer::getMemBufferCopy(
255+
OffloadBinary::write(Image), Image.Image->getBufferIdentifier());
256256
auto NewBinaryOrErr = OffloadBinary::create(*Binary);
257257
if (!NewBinaryOrErr)
258258
return NewBinaryOrErr.takeError();
@@ -1358,6 +1358,135 @@ static Expected<StringRef> linkDevice(ArrayRef<StringRef> InputFiles,
13581358
return *DeviceLinkedFile;
13591359
}
13601360

1361+
static bool isStaticArchiveFile(const StringRef Filename) {
1362+
if (!llvm::sys::path::has_extension(Filename))
1363+
// Any file with no extension should not be considered an Archive.
1364+
return false;
1365+
llvm::file_magic Magic;
1366+
llvm::identify_magic(Filename, Magic);
1367+
// Only archive files are to be considered.
1368+
// TODO: .lib check to be added
1369+
return (Magic == llvm::file_magic::archive);
1370+
}
1371+
1372+
static Expected<StringRef> listSection(StringRef Filename,
1373+
const ArgList &Args) {
1374+
Expected<std::string> OffloadBundlerPath = findProgram(
1375+
"clang-offload-bundler", {getMainExecutable("clang-offload-bundler")});
1376+
if (!OffloadBundlerPath)
1377+
return OffloadBundlerPath.takeError();
1378+
BumpPtrAllocator Alloc;
1379+
StringSaver Saver(Alloc);
1380+
1381+
SmallVector<StringRef, 8> CmdArgs;
1382+
CmdArgs.push_back(*OffloadBundlerPath);
1383+
bool IsArchive = isStaticArchiveFile(Filename);
1384+
CmdArgs.push_back(IsArchive ? "-type=aoo" : "-type=o");
1385+
CmdArgs.push_back(Saver.save("-input=" + Filename));
1386+
CmdArgs.push_back("-list");
1387+
auto Output = createOutputFile("bundled-targets", "list");
1388+
if (!Output)
1389+
return Output.takeError();
1390+
SmallVector<std::optional<StringRef>> Redirects{std::nullopt, *Output,
1391+
std::nullopt};
1392+
int ErrCode = llvm::sys::ExecuteAndWait(*OffloadBundlerPath, CmdArgs,
1393+
std::nullopt, Redirects);
1394+
if (ErrCode != 0)
1395+
return createStringError(inconvertibleErrorCode(),
1396+
"Failed to list targets");
1397+
return *Output;
1398+
}
1399+
1400+
// This routine is used to run the clang-offload-bundler tool and unbundle
1401+
// device inputs that have been created with an older compiler where the
1402+
// device object is bundled into a host object.
1403+
static Expected<StringRef> unbundle(StringRef Filename, const ArgList &Args,
1404+
llvm::Triple Triple) {
1405+
Expected<std::string> OffloadBundlerPath = findProgram(
1406+
"clang-offload-bundler", {getMainExecutable("clang-offload-bundler")});
1407+
if (!OffloadBundlerPath)
1408+
return OffloadBundlerPath.takeError();
1409+
1410+
// Create a new file to write the unbundled file to.
1411+
auto TempFileOrErr =
1412+
createOutputFile(sys::path::filename(ExecutableName), "ir");
1413+
if (!TempFileOrErr)
1414+
return TempFileOrErr.takeError();
1415+
1416+
BumpPtrAllocator Alloc;
1417+
StringSaver Saver(Alloc);
1418+
1419+
SmallVector<StringRef, 8> CmdArgs;
1420+
CmdArgs.push_back(*OffloadBundlerPath);
1421+
bool IsArchive = isStaticArchiveFile(Filename);
1422+
CmdArgs.push_back(IsArchive ? "-type=aoo" : "-type=o");
1423+
auto *Target = Args.MakeArgString(Twine("-targets=sycl-") + Triple.str());
1424+
CmdArgs.push_back(Target);
1425+
CmdArgs.push_back(Saver.save("-input=" + Filename));
1426+
CmdArgs.push_back(Saver.save("-output=" + *TempFileOrErr));
1427+
CmdArgs.push_back("-unbundle");
1428+
CmdArgs.push_back("-allow-missing-bundles");
1429+
if (Error Err = executeCommands(*OffloadBundlerPath, CmdArgs))
1430+
return std::move(Err);
1431+
return *TempFileOrErr;
1432+
}
1433+
1434+
Error extractBundledObjects(StringRef Filename, const ArgList &Args,
1435+
SmallVector<OffloadFile> &Binaries) {
1436+
auto List = listSection(Filename, Args);
1437+
if (!List)
1438+
return List.takeError();
1439+
SmallVector<StringRef> TriplesInFile;
1440+
llvm::ErrorOr<std::unique_ptr<MemoryBuffer>> TripleList =
1441+
llvm::MemoryBuffer::getFileOrSTDIN(*List, /*isText=*/true);
1442+
if (std::error_code EC = TripleList.getError())
1443+
return createFileError(*List, EC);
1444+
(*TripleList)
1445+
->getBuffer()
1446+
.split(TriplesInFile, '\n', /*MaxSplit=*/-1, /*KeepEmpty=*/false);
1447+
for (StringRef TripleStr : TriplesInFile) {
1448+
StringRef SYCLPrefix = "sycl-";
1449+
if (!TripleStr.starts_with(SYCLPrefix))
1450+
continue;
1451+
llvm::Triple Triple(TripleStr.substr(SYCLPrefix.size()));
1452+
auto UnbundledFile = unbundle(Filename, Args, Triple);
1453+
if (!UnbundledFile)
1454+
return UnbundledFile.takeError();
1455+
if (*UnbundledFile == Filename)
1456+
continue;
1457+
1458+
SmallVector<StringRef> ObjectFilePaths;
1459+
if (sycl::isStaticArchiveFile(Filename)) {
1460+
llvm::ErrorOr<std::unique_ptr<MemoryBuffer>> ObjList =
1461+
llvm::MemoryBuffer::getFileOrSTDIN(*UnbundledFile, /*isText=*/true);
1462+
if (std::error_code EC = ObjList.getError())
1463+
return createFileError(*UnbundledFile, EC);
1464+
(*ObjList)->getBuffer().split(ObjectFilePaths, '\n', /*MaxSplit=*/-1,
1465+
/*KeepEmpty=*/false);
1466+
} else {
1467+
ObjectFilePaths.push_back(*UnbundledFile);
1468+
}
1469+
for (StringRef ObjectFilePath : ObjectFilePaths) {
1470+
llvm::file_magic Magic;
1471+
llvm::identify_magic(ObjectFilePath, Magic);
1472+
if (Magic == file_magic::spirv_object)
1473+
return createStringError(
1474+
"SPIR-V fat objects must be generated with --offload-new-driver");
1475+
auto Arg = Args.MakeArgString(
1476+
"sycl-" +
1477+
(Triple.isSPIROrSPIRV() ? Triple.str() + "-" : Triple.str()) + "=" +
1478+
ObjectFilePath);
1479+
auto Binary = getInputBitcodeLibrary(Arg);
1480+
1481+
if (!Binary)
1482+
return Binary.takeError();
1483+
1484+
Binaries.push_back(std::move(*Binary));
1485+
}
1486+
}
1487+
return Error::success();
1488+
}
1489+
13611490
} // namespace sycl
13621491

13631492
namespace generic {
@@ -2634,8 +2763,14 @@ getDeviceInput(const ArgList &Args) {
26342763
if (identify_magic(Buffer.getBuffer()) == file_magic::elf_shared_object)
26352764
continue;
26362765
SmallVector<OffloadFile> Binaries;
2766+
size_t OldSize = Binaries.size();
26372767
if (Error Err = extractOffloadBinaries(Buffer, Binaries))
26382768
return std::move(Err);
2769+
if (Binaries.size() == OldSize) {
2770+
if (Error Err = sycl::extractBundledObjects(*Filename, Args, Binaries))
2771+
return std::move(Err);
2772+
}
2773+
26392774
for (auto &OffloadFile : Binaries) {
26402775
if (identify_magic(Buffer.getBuffer()) == file_magic::archive &&
26412776
!WholeArchive)

devops/dependencies-igc-dev.json

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
11
{
22
"linux": {
33
"igc_dev": {
4-
"github_tag": "igc-dev-427a492",
5-
"version": "427a492",
6-
"updated_at": "2024-08-27T03:48:42Z",
7-
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1857739280/zip",
4+
"github_tag": "igc-dev-fd82ad7",
5+
"version": "fd82ad7",
6+
"updated_at": "2024-09-12T13:46:06Z",
7+
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1924991216/zip",
88
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
99
}
1010
}

devops/dependencies.json

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -19,9 +19,9 @@
1919
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
2020
},
2121
"level_zero": {
22-
"github_tag": "v1.17.39",
23-
"version": "v1.17.39",
24-
"url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.17.39",
22+
"github_tag": "v1.17.42",
23+
"version": "v1.17.42",
24+
"url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.17.42",
2525
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
2626
},
2727
"tbb": {

llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -34,9 +34,7 @@ using EntryPointSet = SetVector<Function *>;
3434

3535
PropSetRegTy computeModuleProperties(const Module &M,
3636
const EntryPointSet &EntryPoints,
37-
const GlobalBinImageProps &GlobProps,
38-
bool SpecConstsMet,
39-
bool IsSpecConstantDefault);
37+
const GlobalBinImageProps &GlobProps);
4038

4139
std::string computeModuleSymbolTable(const Module &M,
4240
const EntryPointSet &EntryPoints);

0 commit comments

Comments
 (0)