Skip to content

Commit 3a66fd7

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (15 commits)
2 parents 809e652 + 8825863 commit 3a66fd7

File tree

48 files changed

+316
-1530
lines changed

Some content is hidden

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

48 files changed

+316
-1530
lines changed

‎clang/include/clang/Basic/DiagnosticDriverKinds.td

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -522,6 +522,9 @@ def warn_drv_no_floating_point_registers: Warning<
522522
InGroup<UnsupportedABI>;
523523
def warn_ignoring_ftabstop_value : Warning<
524524
"ignoring invalid -ftabstop value '%0', using default value %1">;
525+
def warn_ignoring_value_using_default : Warning<
526+
"ignoring invalid '%0' value '%1', using default value '%2'">,
527+
InGroup<UnusedCommandLineArgument>;
525528
def warn_drv_overriding_option : Warning<
526529
"overriding '%0' option with '%1'">,
527530
InGroup<DiagGroup<"overriding-option">>;

‎clang/include/clang/Driver/Options.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6841,7 +6841,7 @@ def fintelfpga : Flag<["-"], "fintelfpga">,
68416841
MarshallingInfoFlag<LangOpts<"IntelFPGA">>,
68426842
HelpText<"Perform ahead-of-time compilation for FPGA">;
68436843
def fsycl_device_only : Flag<["-"], "fsycl-device-only">,
6844-
Alias<offload_device_only>, HelpText<"Compile SYCL kernels for device">;
6844+
HelpText<"Compile SYCL kernels for device">;
68456845
def fsycl_embed_ir : Flag<["-"], "fsycl-embed-ir">,
68466846
HelpText<"Embed LLVM IR for runtime kernel fusion">;
68476847
defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-mem",

‎clang/lib/Driver/Driver.cpp

Lines changed: 20 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -854,19 +854,6 @@ static bool addSYCLDefaultTriple(Compilation &C,
854854
return true;
855855
}
856856

857-
// Special function that checks if -fsycl-device-only was passed on the
858-
// command line. -fsycl-device-only is an alias of --offload-device-only.
859-
static bool hasSYCLDeviceOnly(const ArgList &Args) {
860-
if (const Arg *SYCLDeviceOnlyArg =
861-
Args.getLastArg(options::OPT_offload_device_only)) {
862-
while (SYCLDeviceOnlyArg->getAlias())
863-
SYCLDeviceOnlyArg = SYCLDeviceOnlyArg->getAlias();
864-
if (SYCLDeviceOnlyArg->getSpelling().contains("sycl-device-only"))
865-
return true;
866-
}
867-
return false;
868-
}
869-
870857
void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
871858
InputList &Inputs) {
872859

@@ -1089,7 +1076,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
10891076
bool HasValidSYCLRuntime =
10901077
C.getInputArgs().hasFlag(options::OPT_fsycl, options::OPT_fno_sycl,
10911078
false) ||
1092-
hasSYCLDeviceOnly(C.getInputArgs());
1079+
C.getInputArgs().hasArg(options::OPT_fsycl_device_only);
10931080

10941081
Arg *SYCLfpga = C.getInputArgs().getLastArg(options::OPT_fintelfpga);
10951082

@@ -1172,6 +1159,19 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
11721159
C.getInputArgs().getLastArg(options::OPT_fsycl_range_rounding_EQ);
11731160
checkSingleArgValidity(RangeRoundingPreference, {"disable", "force", "on"});
11741161

1162+
// Evaluation of -fsycl-device-obj is slightly different, we will emit
1163+
// a warning and inform the user of the default behavior used.
1164+
// TODO: General usage of this option is to check for 'spirv' and fallthrough
1165+
// to using llvmir. This can be improved to be more obvious in usage.
1166+
if (Arg *DeviceObj = C.getInputArgs().getLastArgNoClaim(
1167+
options::OPT_fsycl_device_obj_EQ)) {
1168+
StringRef ArgValue(DeviceObj->getValue());
1169+
SmallVector<StringRef, 2> DeviceObjValues = {"spirv", "llvmir"};
1170+
if (llvm::find(DeviceObjValues, ArgValue) == DeviceObjValues.end())
1171+
Diag(clang::diag::warn_ignoring_value_using_default)
1172+
<< DeviceObj->getSpelling().split('=').first << ArgValue << "llvmir";
1173+
}
1174+
11751175
Arg *SYCLForceTarget =
11761176
getArgRequiringSYCLRuntime(options::OPT_fsycl_force_target_EQ);
11771177
if (SYCLForceTarget) {
@@ -1758,12 +1758,13 @@ Compilation *Driver::BuildCompilation(ArrayRef<const char *> ArgList) {
17581758
if (Args.getLastArg(options::OPT_fsycl_dump_device_code_EQ))
17591759
DumpDeviceCode = true;
17601760

1761-
if (const Arg *A = Args.getLastArg(options::OPT_offload_host_only,
1762-
options::OPT_offload_device_only,
1763-
options::OPT_offload_host_device)) {
1761+
if (const Arg *A = Args.getLastArg(
1762+
options::OPT_offload_host_only, options::OPT_offload_device_only,
1763+
options::OPT_fsycl_device_only, options::OPT_offload_host_device)) {
17641764
if (A->getOption().matches(options::OPT_offload_host_only))
17651765
Offload = OffloadHost;
1766-
else if (A->getOption().matches(options::OPT_offload_device_only))
1766+
else if (A->getOption().matches(options::OPT_offload_device_only) ||
1767+
A->getOption().matches(options::OPT_fsycl_device_only))
17671768
Offload = OffloadDevice;
17681769
else
17691770
Offload = OffloadHostDevice;
@@ -3139,7 +3140,7 @@ void Driver::BuildInputs(const ToolChain &TC, DerivedArgList &Args,
31393140
Arg *InputTypeArg = nullptr;
31403141
bool IsSYCL =
31413142
Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false) ||
3142-
hasSYCLDeviceOnly(Args);
3143+
Args.hasArg(options::OPT_fsycl_device_only);
31433144

31443145
// The last /TC or /TP option sets the input type to C or C++ globally.
31453146
if (Arg *TCTP = Args.getLastArgNoClaim(options::OPT__SLASH_TC,

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10828,7 +10828,7 @@ static void addArgs(ArgStringList &DstArgs, const llvm::opt::ArgList &Alloc,
1082810828
}
1082910829
}
1083010830

10831-
static bool supportDynamicLinking(const llvm::opt::ArgList &TCArgs) {
10831+
static bool allowDeviceDependencies(const llvm::opt::ArgList &TCArgs) {
1083210832
if (TCArgs.hasFlag(options::OPT_fsycl_allow_device_dependencies,
1083310833
options::OPT_fno_sycl_allow_device_dependencies, false))
1083410834
return true;
@@ -10862,8 +10862,8 @@ static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
1086210862
options::OPT_fsycl_esimd_force_stateless_mem, false))
1086310863
addArgs(PostLinkArgs, TCArgs, {"-lower-esimd-force-stateless-mem=false"});
1086410864

10865-
if (supportDynamicLinking(TCArgs))
10866-
addArgs(PostLinkArgs, TCArgs, {"-support-dynamic-linking"});
10865+
if (allowDeviceDependencies(TCArgs))
10866+
addArgs(PostLinkArgs, TCArgs, {"-allow-device-image-dependencies"});
1086710867
}
1086810868

1086910869
// On Intel targets we don't need non-kernel functions as entry points,
@@ -10880,7 +10880,7 @@ static bool shouldEmitOnlyKernelsAsEntryPoints(const ToolChain &TC,
1088010880
return true;
1088110881
// When supporting dynamic linking, non-kernels in a device image can be
1088210882
// called.
10883-
if (supportDynamicLinking(TCArgs))
10883+
if (allowDeviceDependencies(TCArgs))
1088410884
return false;
1088510885
if (Triple.isNVPTX() || Triple.isAMDGPU())
1088610886
return false;

‎clang/test/Driver/sycl-device.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,9 +3,11 @@
33
// RUN: | FileCheck -check-prefix=CHECK-DEFAULT %s
44
// CHECK-DEFAULT-NOT: "-fsycl-is-device"
55

6-
/// Check "-fsycl-is-device" is passed when compiling for device:
6+
/// Check "-fsycl-is-device" is passed when compiling for device, including when --config is used:
77
// RUN: %clang -### -fsycl-device-only %s 2>&1 \
88
// RUN: | FileCheck -check-prefix=CHECK-SYCL-DEV %s
9+
// RUN: %clang -### --config=%S/Inputs/empty.cfg -fsycl-device-only %s 2>&1 \
10+
// RUN: | FileCheck -check-prefix=CHECK-SYCL-DEV %s
911
// CHECK-SYCL-DEV: "-fsycl-is-device"{{.*}} "-internal-isystem" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl" "-internal-isystem" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include"
1012

1113
/// Check that "-Wno-sycl-strict" is set on compiler invocation with "-fsycl"

‎clang/test/Driver/sycl-offload-old-model.cpp

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -177,16 +177,16 @@
177177
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_OPT_NO_PASS %s
178178
// CHECK_SYCL_POST_LINK_OPT_NO_PASS-NOT: sycl-post-link{{.*}}emit-only-kernels-as-entry-points
179179

180-
/// Check selective passing of -support-dynamic-linking to sycl-post-link tool
181-
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_fpga -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SADD_PASS %s
182-
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SADD_PASS %s
183-
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fno-sycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SADD_NO_PASS %s
184-
// RUN: %clang_cl -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_fpga -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SADD_PASS %s
185-
// RUN: %clang_cl -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SADD_PASS %s
186-
// RUN: %clang_cl -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fno-sycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_SADD_NO_PASS %s
187-
188-
// CHECK_SYCL_POST_LINK_SADD_PASS: sycl-post-link{{.*}}support-dynamic-linking
189-
// CHECK_SYCL_POST_LINK_SADD_NO_PASS-NOT: sycl-post-link{{.*}}support-dynamic-linking
180+
/// Check selective passing of -allow-device-image-dependencies to sycl-post-link tool
181+
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_fpga -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_ADID_PASS %s
182+
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_ADID_PASS %s
183+
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fno-sycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_ADID_NO_PASS %s
184+
// RUN: %clang_cl -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_fpga -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_ADID_PASS %s
185+
// RUN: %clang_cl -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fsycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_ADID_PASS %s
186+
// RUN: %clang_cl -### -fsycl --no-offload-new-driver -fsycl-targets=spir64_gen -fno-sycl-allow-device-dependencies %s 2>&1 | FileCheck -check-prefix=CHECK_SYCL_POST_LINK_ADID_NO_PASS %s
187+
188+
// CHECK_SYCL_POST_LINK_ADID_PASS: sycl-post-link{{.*}}allow-device-image-dependencies
189+
// CHECK_SYCL_POST_LINK_ADID_NO_PASS-NOT: sycl-post-link{{.*}}allow-device-image-dependencies
190190

191191
/// Check for correct handling of -fsycl-fp64-conv-emu option for different targets
192192
// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --no-offload-new-driver -fsycl-targets=spir64 -fsycl-fp64-conv-emu %s 2>&1 | FileCheck -check-prefix=CHECK_WARNING %s

‎clang/test/Driver/sycl-offload.c

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -69,6 +69,13 @@
6969
// RUN: | FileCheck -check-prefix=CHK-SYCL-TARGET %s
7070
// CHK-SYCL-TARGET-NOT: error: SYCL target is invalid
7171

72+
/// -fsycl-device-obj argument check
73+
// RUN: %clang -### -fsycl-device-obj=test -fsycl %s 2>&1 \
74+
// RUN: | FileCheck -check-prefix=DEVICE_OBJ_WARN %s
75+
// RUN: %clang_cl -### -fsycl-device-obj=test -fsycl %s 2>&1 \
76+
// RUN: | FileCheck -check-prefix=DEVICE_OBJ_WARN %s
77+
// DEVICE_OBJ_WARN: warning: ignoring invalid '-fsycl-device-obj' value 'test', using default value 'llvmir' [-Wunused-command-line-argument]
78+
7279
/// ###########################################################################
7380

7481
/// Check warning for duplicate offloading targets.

‎clang/tools/clang-offload-extract/ClangOffloadExtract.cpp

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -62,6 +62,11 @@ to the stem:
6262
)"),
6363
cl::cat(ClangOffloadExtractCategory));
6464

65+
static cl::opt<bool>
66+
Quiet("q", cl::init(false),
67+
cl::desc(R"(Do not print the names of generated files)"),
68+
cl::cat(ClangOffloadExtractCategory));
69+
6570
// Create an alias for the deprecated option, so legacy use still works
6671
static cl::alias FileNameStemAlias(
6772
"output", cl::desc("Deprecated option, replaced by option '--stem'"),
@@ -262,8 +267,10 @@ linked fat binary, and store them in separate files.
262267
OffloadName.erase(0, OffloadPrefix.length());
263268

264269
// Tell user that we are saving an image.
265-
outs() << "Section '" + OffloadName + "': Image " << ImgCnt++
266-
<< "'-> File '" + FileName + "'\n";
270+
if (!Quiet) {
271+
outs() << "Section '" + OffloadName + "': Image " << ImgCnt++
272+
<< "'-> File '" + FileName + "'\n";
273+
}
267274

268275
// Write image data to the output
269276
std::error_code EC;

‎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-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",
4+
"github_tag": "igc-dev-a20debc",
5+
"version": "a20debc",
6+
"updated_at": "2024-09-15T13:44:38Z",
7+
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/1934718090/zip",
88
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
99
}
1010
}

‎llvm-spirv/lib/SPIRV/SPIRVReader.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3044,12 +3044,15 @@ Function *SPIRVToLLVM::transFunction(SPIRVFunction *BF) {
30443044
auto BA = BF->getArgument(I->getArgNo());
30453045
mapValue(BA, &(*I));
30463046
setName(&(*I), BA);
3047+
AttributeMask IllegalAttrs = AttributeFuncs::typeIncompatible(I->getType());
30473048
BA->foreachAttr([&](SPIRVFuncParamAttrKind Kind) {
30483049
// Skip this function parameter attribute as it will translated among
30493050
// OpenCL metadata
30503051
if (Kind == FunctionParameterAttributeRuntimeAlignedINTEL)
30513052
return;
30523053
Attribute::AttrKind LLVMKind = SPIRSPIRVFuncParamAttrMap::rmap(Kind);
3054+
if (IllegalAttrs.contains(LLVMKind))
3055+
return;
30533056
Type *AttrTy = nullptr;
30543057
switch (LLVMKind) {
30553058
case Attribute::AttrKind::ByVal:

‎llvm-spirv/lib/SPIRV/SPIRVWriter.cpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -917,10 +917,7 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {
917917
if (Attrs.hasParamAttr(ArgNo, Attribute::ReadOnly))
918918
BA->addAttr(FunctionParameterAttributeNoWrite);
919919
if (Attrs.hasParamAttr(ArgNo, Attribute::ReadNone))
920-
// TODO: intel/llvm customization
921-
// see https://github.com/intel/llvm/issues/7592
922-
// Need to return FunctionParameterAttributeNoReadWrite
923-
BA->addAttr(FunctionParameterAttributeNoWrite);
920+
BA->addAttr(FunctionParameterAttributeNoReadWrite);
924921
if (Attrs.hasParamAttr(ArgNo, Attribute::ZExt))
925922
BA->addAttr(FunctionParameterAttributeZext);
926923
if (Attrs.hasParamAttr(ArgNo, Attribute::SExt))

‎llvm-spirv/test/image.ll

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,3 @@
1-
; XFAIL: *
2-
31
; RUN: llvm-as %s -o %t.bc
42
; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
53
; RUN: llvm-spirv %t.bc -o %t.spv

‎llvm-spirv/test/transcoding/builtin_function_readnone_attr.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,9 @@
99
; CHECK-SPIRV: Name [[#B:]] "b"
1010
; CHECK-SPIRV: Decorate [[#A]] FuncParamAttr 5
1111
; CHECK-SPIRV: Decorate [[#A]] FuncParamAttr 6
12-
; CHECK-SPIRV: Decorate [[#B]] FuncParamAttr 5
12+
; CHECK-SPIRV: Decorate [[#B]] FuncParamAttr 7
1313

14-
; CHECK-LLVM: {{.*}}void @test_builtin_readnone(ptr nocapture readonly %{{.*}}, ptr nocapture readonly %{{.*}})
14+
; CHECK-LLVM: {{.*}}void @test_builtin_readnone(ptr nocapture readonly %{{.*}}, ptr nocapture readnone %{{.*}})
1515

1616
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
1717
target triple = "spir-unknown-unknown"

‎llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -52,9 +52,9 @@ constexpr char SYCL_SCOPE_NAME[] = "<SYCL>";
5252
constexpr char ESIMD_SCOPE_NAME[] = "<ESIMD>";
5353
constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd";
5454

55-
cl::opt<bool> SupportDynamicLinking{
56-
"support-dynamic-linking",
57-
cl::desc("Generate device images that are suitable for dynamic linking"),
55+
cl::opt<bool> AllowDeviceImageDependencies{
56+
"allow-device-image-dependencies",
57+
cl::desc("Allow dependencies between device images"),
5858
cl::cat(getModuleSplitCategory()), cl::init(false)};
5959

6060
EntryPointsGroupScope selectDeviceCodeGroupScope(const Module &M,
@@ -670,7 +670,7 @@ static bool mustPreserveGV(const GlobalValue &GV) {
670670
// When dynamic linking is supported, we internalize everything that can
671671
// not be imported which also means that there is no point of having it
672672
// visible outside of the current module.
673-
if (SupportDynamicLinking)
673+
if (AllowDeviceImageDependencies)
674674
return canBeImportedFunction(*F);
675675

676676
// Otherwise, we are being even more aggressive: SYCL modules are expected
@@ -715,7 +715,7 @@ void ModuleDesc::cleanup() {
715715

716716
// Callback for internalize can't be a lambda with captures, so we propagate
717717
// necessary information through the module itself.
718-
if (!SupportDynamicLinking)
718+
if (!AllowDeviceImageDependencies)
719719
for (Function *F : EntryPoints.Functions)
720720
F->addFnAttr("sycl-entry-point");
721721

@@ -1391,7 +1391,7 @@ bool canBeImportedFunction(const Function &F) {
13911391
// of user device code (e.g. _Z38__spirv_JointMatrixWorkItemLength...) In
13921392
// order to be safe and not require perfect name analysis just start with this
13931393
// simple check.
1394-
if (!SupportDynamicLinking)
1394+
if (!AllowDeviceImageDependencies)
13951395
return false;
13961396

13971397
// SYCL_EXTERNAL property is not recorded for a declaration

‎llvm/lib/SYCLLowerIR/SpecConstants.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -951,8 +951,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
951951
updatePaddingInLastMDNode(Ctx, SCMetadata, Padding);
952952
}
953953

954+
auto *DefValTy = DefaultValue->getType();
954955
SCMetadata[SymID] = generateSpecConstantMetadata(
955-
M, SymID, SCTy, NextID, /* is native spec constant */ false);
956+
M, SymID, DefValTy, NextID, /* is native spec constant */ false);
956957

957958
++NextID.ID;
958959
NextOffset += Size;

‎llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@ add_llvm_component_library(LLVMSYCLNativeCPUUtils
1313
Core
1414
Support
1515
Passes
16+
SYCLLowerIR
1617
Target
1718
TargetParser
1819
TransformUtils

‎llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,7 @@
1414
#include "llvm/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.h"
1515
#include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
1616
#include "llvm/SYCLLowerIR/RenameKernelSYCLNativeCPU.h"
17+
#include "llvm/SYCLLowerIR/SpecConstants.h"
1718
#include "llvm/SYCLLowerIR/UtilsSYCLNativeCPU.h"
1819
#include "llvm/Support/CommandLine.h"
1920

@@ -60,6 +61,7 @@ static cl::opt<bool>
6061
void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
6162
llvm::ModulePassManager &MPM, ModuleAnalysisManager &MAM,
6263
OptimizationLevel OptLevel) {
64+
MPM.addPass(SpecConstantsPass(SpecConstantsPass::HandlingMode::emulation));
6365
MPM.addPass(ConvertToMuxBuiltinsSYCLNativeCPUPass());
6466
#ifdef NATIVECPU_USE_OCK
6567
MPM.addPass(compiler::utils::TransferKernelMetadataPass());

‎llvm/test/tools/sycl-post-link/emit_imported_symbols.ll

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,12 @@
11
; This test checks that the -emit-imported-symbols option generates a list of imported symbols
22
; Function names were chosen so that no function with a 'inside' in their function name is imported
3-
; Note that -emit-imported-symbols will not emit any imported symbols without -support-dynamic-linking.
3+
; Note that -emit-imported-symbols will not emit any imported symbols without -allow-device-image-dependencies.
44

55
;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
66
; Test with -split=kernel
77
;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
88

9-
; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -emit-imported-symbols -split=kernel -S < %s -o %t_kernel.table
9+
; RUN: sycl-post-link -properties -symbols -allow-device-image-dependencies -emit-imported-symbols -split=kernel -S < %s -o %t_kernel.table
1010

1111
; RUN: FileCheck %s -input-file=%t_kernel_0.sym --check-prefixes CHECK-KERNEL-SYM-0
1212
; RUN: FileCheck %s -input-file=%t_kernel_1.sym --check-prefixes CHECK-KERNEL-SYM-1
@@ -41,11 +41,11 @@
4141
; Test with -split=source
4242
;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
4343

44-
; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -emit-imported-symbols -split=source -S < %s -o %t_source.table
44+
; RUN: sycl-post-link -properties -symbols -allow-device-image-dependencies -emit-imported-symbols -split=source -S < %s -o %t_source.table
4545
; RUN: FileCheck %s -input-file=%t_source_0.sym --check-prefixes CHECK-SOURCE-SYM-0
4646
; RUN: FileCheck %s -input-file=%t_source_0.prop --check-prefixes CHECK-SOURCE-IMPORTED-SYM-0
4747

48-
; RUN: sycl-post-link -properties -symbols -support-dynamic-linking -emit-imported-symbols -split=source -S < %s -o %t_source.table -O0
48+
; RUN: sycl-post-link -properties -symbols -allow-device-image-dependencies -emit-imported-symbols -split=source -S < %s -o %t_source.table -O0
4949
; RUN: FileCheck %s -input-file=%t_source_0.sym --check-prefixes CHECK-SOURCE-SYM-0
5050
; RUN: FileCheck %s -input-file=%t_source_0.prop --check-prefixes CHECK-SOURCE-IMPORTED-SYM-0
5151

0 commit comments

Comments
 (0)