Skip to content

Commit 5fbd5b0

Browse files
joppermjopperm
committed
Merge remote-tracking branch 'upstream/sycl' into rtc-includes
2 parents f008e0a + cce5dac commit 5fbd5b0

File tree

13 files changed

+331
-75
lines changed

13 files changed

+331
-75
lines changed

‎.github/workflows/coverity.yml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ jobs:
5050
--cmake-opt="-DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD=SPIRV"
5151
5252
- name: Build with coverity
53-
run: $GITHUB_WORKSPACE/cov-analysis-linux64-*/bin/cov-build --dir cov-int cmake --build $GITHUB_WORKSPACE/build
53+
run: $GITHUB_WORKSPACE/cov-analysis-linux64-*/bin/cov-build --dir cov-int cmake --build $GITHUB_WORKSPACE/build --target sycl-toolchain
5454

5555
- name: Compress results
5656
run: tar -I pigz -cf intel_llvm.tgz cov-int

‎.github/workflows/sycl-windows-build.yml

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -104,6 +104,7 @@ jobs:
104104
echo "C:\Program Files\Git\usr\bin" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append
105105
echo "CCACHE_DIR=D:\github\_work\cache\${{ inputs.build_cache_suffix }}" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
106106
echo "CCACHE_MAXSIZE=10G" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
107+
echo "LIT_OPTS='-j$env:NUMBER_OF_PROCESSORS $LIT_OPTS'" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
107108
- name: Register cleanup after job is finished
108109
uses: ./devops/actions/cleanup
109110
- uses: ./devops/actions/cached_checkout

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10769,6 +10769,8 @@ static void getNonTripleBasedSPIRVTransOpts(Compilation &C,
1076910769
// to user supplied options.
1077010770
// NOTE: Any changes made here should be reflected in the similarly named
1077110771
// function in clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp.
10772+
// NOTE2: JIT related changes made here should be reflected in 'translatorOpts'
10773+
// from sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp.
1077210774
static void getTripleBasedSPIRVTransOpts(Compilation &C,
1077310775
const llvm::opt::ArgList &TCArgs,
1077410776
llvm::Triple Triple,

‎sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

Lines changed: 62 additions & 59 deletions
Original file line numberDiff line numberDiff line change
@@ -660,79 +660,82 @@ jit_compiler::performPostLink(std::unique_ptr<llvm::Module> Module,
660660
// TODO: Call `verifyNoCrossModuleDeviceGlobalUsage` if device globals shall
661661
// be processed.
662662

663-
// TODO: This allocation assumes that there are no further splits required,
664-
// i.e. there are no mixed SYCL/ESIMD modules.
665-
RTCBundleInfo BundleInfo;
666-
BundleInfo.DevImgInfos = DynArray<RTCDevImgInfo>{Splitter->remainingSplits()};
663+
SmallVector<RTCDevImgInfo> DevImgInfoVec;
667664
SmallVector<std::unique_ptr<llvm::Module>> Modules;
668665

669-
auto *DevImgInfoIt = BundleInfo.DevImgInfos.begin();
670-
while (Splitter->hasMoreSplits()) {
671-
assert(DevImgInfoIt != BundleInfo.DevImgInfos.end());
666+
// TODO: The following logic is missing the ability to link ESIMD and SYCL
667+
// modules back together, which would be requested via
668+
// `-fno-sycl-device-code-split-esimd` as a prerequisite for compiling
669+
// `invoke_simd` code.
672670

671+
while (Splitter->hasMoreSplits()) {
673672
ModuleDesc MDesc = Splitter->nextSplit();
674-
RTCDevImgInfo &DevImgInfo = *DevImgInfoIt++;
675673

676674
// TODO: Call `MDesc.fixupLinkageOfDirectInvokeSimdTargets()` when
677675
// `invoke_simd` is supported.
678676

679677
SmallVector<ModuleDesc, 2> ESIMDSplits =
680678
splitByESIMD(std::move(MDesc), EmitOnlyKernelsAsEntryPoints);
681-
assert(!ESIMDSplits.empty());
682-
if (ESIMDSplits.size() > 1) {
683-
return createStringError("Mixing SYCL and ESIMD code is unsupported");
684-
}
685-
MDesc = std::move(ESIMDSplits.front());
686-
687-
if (MDesc.isESIMD()) {
688-
// `sycl-post-link` has a `-lower-esimd` option, but there's no clang
689-
// driver option to influence it. Rather, the driver sets it
690-
// unconditionally in the multi-file output mode, which we are mimicking
691-
// here.
692-
lowerEsimdConstructs(MDesc, PerformOpts);
693-
}
694-
695-
MDesc.saveSplitInformationAsMetadata();
696-
697-
DevImgInfo.SymbolTable = FrozenSymbolTable{MDesc.entries().size()};
698-
transform(MDesc.entries(), DevImgInfo.SymbolTable.begin(),
699-
[](Function *F) { return F->getName(); });
700-
701-
// TODO: Determine what is requested.
702-
GlobalBinImageProps PropReq{
703-
/*EmitKernelParamInfo=*/true, /*EmitProgramMetadata=*/true,
704-
/*EmitExportedSymbols=*/true, /*EmitImportedSymbols=*/true,
705-
/*DeviceGlobals=*/false};
706-
PropertySetRegistry Properties =
707-
computeModuleProperties(MDesc.getModule(), MDesc.entries(), PropReq);
708-
// TODO: Manually add `compile_target` property as in
709-
// `saveModuleProperties`?
710-
const auto &PropertySets = Properties.getPropSets();
711-
712-
DevImgInfo.Properties = FrozenPropertyRegistry{PropertySets.size()};
713-
for (auto [KV, FrozenPropSet] :
714-
zip_equal(PropertySets, DevImgInfo.Properties)) {
715-
const auto &PropertySetName = KV.first;
716-
const auto &PropertySet = KV.second;
717-
FrozenPropSet =
718-
FrozenPropertySet{PropertySetName.str(), PropertySet.size()};
719-
for (auto [KV2, FrozenProp] :
720-
zip_equal(PropertySet, FrozenPropSet.Values)) {
721-
const auto &PropertyName = KV2.first;
722-
const auto &PropertyValue = KV2.second;
723-
FrozenProp =
724-
PropertyValue.getType() == PropertyValue::Type::UINT32
725-
? FrozenPropertyValue{PropertyName.str(),
726-
PropertyValue.asUint32()}
727-
: FrozenPropertyValue{PropertyName.str(),
728-
PropertyValue.asRawByteArray(),
729-
PropertyValue.getRawByteArraySize()};
679+
for (auto &ES : ESIMDSplits) {
680+
MDesc = std::move(ES);
681+
682+
if (MDesc.isESIMD()) {
683+
// `sycl-post-link` has a `-lower-esimd` option, but there's no clang
684+
// driver option to influence it. Rather, the driver sets it
685+
// unconditionally in the multi-file output mode, which we are mimicking
686+
// here.
687+
lowerEsimdConstructs(MDesc, PerformOpts);
730688
}
731-
};
732689

733-
Modules.push_back(MDesc.releaseModulePtr());
690+
MDesc.saveSplitInformationAsMetadata();
691+
692+
RTCDevImgInfo &DevImgInfo = DevImgInfoVec.emplace_back();
693+
DevImgInfo.SymbolTable = FrozenSymbolTable{MDesc.entries().size()};
694+
transform(MDesc.entries(), DevImgInfo.SymbolTable.begin(),
695+
[](Function *F) { return F->getName(); });
696+
697+
// TODO: Determine what is requested.
698+
GlobalBinImageProps PropReq{
699+
/*EmitKernelParamInfo=*/true, /*EmitProgramMetadata=*/true,
700+
/*EmitExportedSymbols=*/true, /*EmitImportedSymbols=*/true,
701+
/*DeviceGlobals=*/false};
702+
PropertySetRegistry Properties =
703+
computeModuleProperties(MDesc.getModule(), MDesc.entries(), PropReq);
704+
// TODO: Manually add `compile_target` property as in
705+
// `saveModuleProperties`?
706+
const auto &PropertySets = Properties.getPropSets();
707+
708+
DevImgInfo.Properties = FrozenPropertyRegistry{PropertySets.size()};
709+
for (auto [KV, FrozenPropSet] :
710+
zip_equal(PropertySets, DevImgInfo.Properties)) {
711+
const auto &PropertySetName = KV.first;
712+
const auto &PropertySet = KV.second;
713+
FrozenPropSet =
714+
FrozenPropertySet{PropertySetName.str(), PropertySet.size()};
715+
for (auto [KV2, FrozenProp] :
716+
zip_equal(PropertySet, FrozenPropSet.Values)) {
717+
const auto &PropertyName = KV2.first;
718+
const auto &PropertyValue = KV2.second;
719+
FrozenProp =
720+
PropertyValue.getType() == PropertyValue::Type::UINT32
721+
? FrozenPropertyValue{PropertyName.str(),
722+
PropertyValue.asUint32()}
723+
: FrozenPropertyValue{PropertyName.str(),
724+
PropertyValue.asRawByteArray(),
725+
PropertyValue.getRawByteArraySize()};
726+
}
727+
};
728+
729+
Modules.push_back(MDesc.releaseModulePtr());
730+
}
734731
}
735732

733+
assert(DevImgInfoVec.size() == Modules.size());
734+
RTCBundleInfo BundleInfo;
735+
BundleInfo.DevImgInfos = DynArray<RTCDevImgInfo>{DevImgInfoVec.size()};
736+
std::move(DevImgInfoVec.begin(), DevImgInfoVec.end(),
737+
BundleInfo.DevImgInfos.begin());
738+
736739
return PostLinkResult{std::move(BundleInfo), std::move(Modules)};
737740
}
738741

‎sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp

Lines changed: 62 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -23,29 +23,78 @@ using namespace jit_compiler::translation;
2323
using namespace llvm;
2424

2525
SPIRV::TranslatorOpts &SPIRVLLVMTranslator::translatorOpts() {
26-
static auto Opts = []() -> SPIRV::TranslatorOpts {
26+
// Keep this in sync with clang/lib/Driver/ToolChains/Clang.cpp
27+
// TODO: consider introducing a config file that both clang and jit-compiler
28+
// could use during options setting.
29+
std::vector<SPIRV::ExtensionID> AllowedExtensions{
30+
SPIRV::ExtensionID::SPV_EXT_shader_atomic_float_add,
31+
SPIRV::ExtensionID::SPV_EXT_shader_atomic_float_min_max,
32+
SPIRV::ExtensionID::SPV_KHR_no_integer_wrap_decoration,
33+
SPIRV::ExtensionID::SPV_KHR_float_controls,
34+
SPIRV::ExtensionID::SPV_KHR_expect_assume,
35+
SPIRV::ExtensionID::SPV_KHR_linkonce_odr,
36+
SPIRV::ExtensionID::SPV_INTEL_subgroups,
37+
SPIRV::ExtensionID::SPV_INTEL_media_block_io,
38+
SPIRV::ExtensionID::SPV_INTEL_device_side_avc_motion_estimation,
39+
SPIRV::ExtensionID::SPV_INTEL_fpga_loop_controls,
40+
SPIRV::ExtensionID::SPV_INTEL_unstructured_loop_controls,
41+
SPIRV::ExtensionID::SPV_INTEL_fpga_reg,
42+
SPIRV::ExtensionID::SPV_INTEL_blocking_pipes,
43+
SPIRV::ExtensionID::SPV_INTEL_function_pointers,
44+
SPIRV::ExtensionID::SPV_INTEL_kernel_attributes,
45+
SPIRV::ExtensionID::SPV_INTEL_io_pipes,
46+
SPIRV::ExtensionID::SPV_INTEL_inline_assembly,
47+
SPIRV::ExtensionID::SPV_INTEL_arbitrary_precision_integers,
48+
SPIRV::ExtensionID::SPV_INTEL_float_controls2,
49+
SPIRV::ExtensionID::SPV_INTEL_vector_compute,
50+
SPIRV::ExtensionID::SPV_INTEL_fast_composite,
51+
SPIRV::ExtensionID::SPV_INTEL_arbitrary_precision_fixed_point,
52+
SPIRV::ExtensionID::SPV_INTEL_arbitrary_precision_floating_point,
53+
SPIRV::ExtensionID::SPV_INTEL_variable_length_array,
54+
SPIRV::ExtensionID::SPV_INTEL_fp_fast_math_mode,
55+
SPIRV::ExtensionID::SPV_INTEL_long_composites,
56+
SPIRV::ExtensionID::SPV_INTEL_arithmetic_fence,
57+
SPIRV::ExtensionID::SPV_INTEL_global_variable_decorations,
58+
SPIRV::ExtensionID::SPV_INTEL_cache_controls,
59+
SPIRV::ExtensionID::SPV_INTEL_fpga_buffer_location,
60+
SPIRV::ExtensionID::SPV_INTEL_fpga_argument_interfaces,
61+
SPIRV::ExtensionID::SPV_INTEL_fpga_invocation_pipelining_attributes,
62+
SPIRV::ExtensionID::SPV_INTEL_fpga_latency_control,
63+
SPIRV::ExtensionID::SPV_KHR_shader_clock,
64+
SPIRV::ExtensionID::SPV_INTEL_bindless_images,
65+
SPIRV::ExtensionID::SPV_INTEL_task_sequence,
66+
SPIRV::ExtensionID::SPV_INTEL_bfloat16_conversion,
67+
SPIRV::ExtensionID::SPV_INTEL_joint_matrix,
68+
SPIRV::ExtensionID::SPV_INTEL_hw_thread_queries,
69+
SPIRV::ExtensionID::SPV_KHR_uniform_group_instructions,
70+
SPIRV::ExtensionID::SPV_INTEL_masked_gather_scatter,
71+
SPIRV::ExtensionID::SPV_INTEL_tensor_float32_conversion,
72+
SPIRV::ExtensionID::SPV_INTEL_optnone,
73+
SPIRV::ExtensionID::SPV_KHR_non_semantic_info,
74+
SPIRV::ExtensionID::SPV_KHR_cooperative_matrix,
75+
SPIRV::ExtensionID::SPV_EXT_shader_atomic_float16_add,
76+
SPIRV::ExtensionID::SPV_INTEL_fp_max_error};
77+
78+
static auto Opts = [&]() -> SPIRV::TranslatorOpts {
2779
// Options for translation between SPIR-V and LLVM IR.
28-
// Set SPIRV-V 1.4 as the maximum version number for now.
80+
// Set SPIRV-V 1.5 as the maximum version number for now.
2981
// Note that some parts of the code depend on the available builtins, e.g.,
3082
// passes/kernel-fusion/Builtins.cpp, so updating the SPIR-V version should
3183
// involve revisiting that code.
32-
SPIRV::TranslatorOpts TransOpt{SPIRV::VersionNumber::SPIRV_1_4};
84+
SPIRV::TranslatorOpts TransOpt{SPIRV::VersionNumber::SPIRV_1_5};
3385
// Enable attachment of kernel arg names as metadata.
3486
TransOpt.enableGenArgNameMD();
3587
// Enable mem2reg.
3688
TransOpt.setMemToRegEnabled(true);
37-
// Enable all extensions.
38-
// TODO: Specifically enable only the
39-
// extensions listed in the KernelInfo.
40-
// FIXME: Because there's no size provided,
41-
// there's currently no obvious way to iterate the
42-
// array of extensions in KernelInfo.
43-
TransOpt.enableAllExtensions();
44-
// TODO: Remove this workaround.
45-
TransOpt.setAllowedToUseExtension(
46-
SPIRV::ExtensionID::SPV_KHR_untyped_pointers, false);
89+
for (auto &Ext : AllowedExtensions)
90+
TransOpt.setAllowedToUseExtension(Ext, true);
4791
TransOpt.setDesiredBIsRepresentation(
4892
SPIRV::BIsRepresentation::SPIRVFriendlyIR);
93+
TransOpt.setDebugInfoEIS(
94+
SPIRV::DebugInfoEIS::NonSemantic_Shader_DebugInfo_200);
95+
const llvm::SmallVector<llvm::StringRef, 4> AllowedIntrinsics = {
96+
"llvm.genx."};
97+
TransOpt.setSPIRVAllowUnknownIntrinsics(AllowedIntrinsics);
4998
// TODO: We need to take care of specialization constants, either by
5099
// instantiating them by the user-supplied value from the SYCL runtime or by
51100
// making sure they are correctly represented in the output of the fusion

‎sycl/ReleaseNotes.md

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,21 @@
1+
# Release notes for an upcoming release (dates TBD)
2+
3+
## New Features
4+
5+
- Added support for ... intel/llvm#pr
6+
7+
## Improvements
8+
9+
- Improved handling of ... intel/llvm#pr
10+
11+
## Bug Fixes
12+
13+
- Fixed ... intel/llvm#pr
14+
15+
## Misc
16+
17+
- Did this and that ... intel/llvm#pr
18+
119
# Release notes Nov'24
220

321
Release notes for commit range

‎sycl/doc/developer/ContributeToDPCPP.md

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,36 @@ For any DPC++-related commit, the `[SYCL]` tag should be present in the
3737
commit message title. To a reasonable extent, additional tags can be used
3838
to signify the component changed, e.g.: `[UR]`, `[CUDA]`, `[Doc]`.
3939

40+
## Release notes
41+
42+
You are encouraged to record your change into
43+
[release notes](https://github.com/intel/llvm/blob/sycl/sycl/ReleaseNotes.md)
44+
under "Release notes for an upcoming release" section.
45+
46+
A change should be noted there when:
47+
48+
- A public interface (API, command line options, env variables, etc.) is being
49+
changed
50+
- A bug is being fixed
51+
- Any change is being made which has an observable behavior (including
52+
performance)
53+
54+
A change should **not** be noted there when:
55+
56+
- It has no functional or performance impact
57+
- It is about our CI infrastructure, testing infrastructure, or tests
58+
59+
There are no strict guidelines on how to structure release notes, but for
60+
consistency it is better to follow the existing structure minimal changes. The
61+
structure we have been using so far is split by change type (i.e. new
62+
features and bug fixes) and then sub-split by component (i.e. compiler,
63+
runtime). Please use past tense when describing your change and leave a
64+
reference to your PR using `intel/llvm#NNNNN` syntax. If you want to reference
65+
an in-tree document (like an extension spec, for example), it must be
66+
referenced using a permalink so that it stays actual even if document is moved.
67+
68+
When in doubt, you can always ask reviewers for help/guidance/suggestions.
69+
4070
## Using \<iostream\>
4171

4272
According to

‎sycl/doc/developer/WorkingOnAReleaseBranch.md

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
# Working on a release branch
2+
3+
A "release branch" is defined as a branch whose name starts with `sycl-rel-`
4+
prefix.
5+
6+
Those branches are intended to indicate stable snapshots of our product so that
7+
our users don't need to guess which nightly build is good enough for their
8+
needs.
9+
10+
Therefore, those branches have higher quality requirements and as such have
11+
different contribution rules intended to preserve their stability.
12+
13+
If you are not familiar with the [general contribution guidelines][contributing]
14+
or the [DPC++ specific contribution guidelines][contributing-to-dpcpp], please
15+
familiarize yourself with those documents first because they also apply to
16+
release branches.
17+
18+
## Extra rules for release branches
19+
20+
### Only cherry-picks are allowed
21+
22+
It is assumed that everything you do on a release branch should also be
23+
repeated on the default `sycl` branch to ensure that it is automatically
24+
included into future releases.
25+
26+
Therefore, when submitting a PR to a release branch, its description should
27+
contain a link to the corresponding PR in the default `sycl` branch.
28+
29+
Note that it is not acceptable to first merge something into a
30+
release branch and then apply it to the default `sycl` branch. The flow goes in
31+
the opposite direction where you first land a patch to the default `sycl` branch
32+
and then backport it to a release branch.
33+
34+
### No new features are allowed
35+
36+
Features are generally more complicated than bug fixes and may require further
37+
bug fixes as well. Considering that release branches are intended to be stable,
38+
no new features are allowed to be added there.
39+
40+
[contributing]: https://github.com/intel/llvm/blob/sycl/CONTRIBUTING.md
41+
[contributing-to-dpcpp]: ./ContributeToDPCPP.md

‎sycl/doc/index.rst

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -67,3 +67,4 @@ Developer Documentation
6767
developer/ABIPolicyGuide
6868
developer/ContributeToDPCPP
6969
developer/KHRExtensions
70+
developer/WorkingOnAReleaseBranch

‎sycl/source/detail/kernel_bundle_impl.hpp

Lines changed: 15 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -495,7 +495,21 @@ class kernel_bundle_impl {
495495
std::vector<ur_device_handle_t> DeviceVec;
496496
DeviceVec.reserve(Devices.size());
497497
for (const auto &SyclDev : Devices) {
498-
ur_device_handle_t Dev = getSyclObjImpl(SyclDev)->getHandleRef();
498+
DeviceImplPtr DevImpl = getSyclObjImpl(SyclDev);
499+
if (!ContextImpl->hasDevice(DevImpl)) {
500+
throw sycl::exception(make_error_code(errc::invalid),
501+
"device not part of kernel_bundle context");
502+
}
503+
if (!DevImpl->extOneapiCanCompile(MLanguage)) {
504+
// This error cannot not be exercised in the current implementation, as
505+
// compatibility with a source language depends on the backend's
506+
// capabilities and all devices in one context share the same backend in
507+
// the current implementation, so this would lead to an error already
508+
// during construction of the source bundle.
509+
throw sycl::exception(make_error_code(errc::invalid),
510+
"device does not support source language");
511+
}
512+
ur_device_handle_t Dev = DevImpl->getHandleRef();
499513
DeviceVec.push_back(Dev);
500514
}
501515

‎sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//==- kernel_compiler_sycl_jit.cpp --- kernel_compiler extension tests -----==//
1+
//==---- kernel_compiler_basic.cpp --- kernel_compiler extension tests -----==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.

0 commit comments

Comments
 (0)