Skip to content

Commit 921c64b

Browse files
committed
Merge branch 'sycl' into work_group_static
2 parents 3b3f327 + 814290d commit 921c64b

File tree

18 files changed

+189
-42
lines changed

18 files changed

+189
-42
lines changed

clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1789,6 +1789,9 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn,
17891789
if (SyclOptReport.HasOptReportInfo(FD)) {
17901790
llvm::OptimizationRemarkEmitter ORE(Fn);
17911791
for (auto ORI : llvm::enumerate(SyclOptReport.GetInfo(FD))) {
1792+
// Temporarily apply arg location to ensure SourceLocToDebugLoc
1793+
// picks up the expected file.
1794+
ApplyDebugLocation TempApplyLoc(*this, ORI.value().KernelArgLoc);
17921795
llvm::DiagnosticLocation DL =
17931796
SourceLocToDebugLoc(ORI.value().KernelArgLoc);
17941797
StringRef NameInDesc = ORI.value().KernelArgDescName;

clang/lib/Driver/ToolChains/SYCL.cpp

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1618,6 +1618,23 @@ static std::vector<OptSpecifier> getUnsupportedOpts(void) {
16181618
return UnsupportedOpts;
16191619
}
16201620

1621+
// Currently supported options by SYCL NativeCPU device compilation
1622+
static inline bool SupportedByNativeCPU(const SYCLToolChain &TC,
1623+
const OptSpecifier &Opt) {
1624+
if (!TC.IsSYCLNativeCPU)
1625+
return false;
1626+
1627+
switch (Opt.getID()) {
1628+
case options::OPT_fcoverage_mapping:
1629+
case options::OPT_fno_coverage_mapping:
1630+
case options::OPT_fprofile_instr_generate:
1631+
case options::OPT_fprofile_instr_generate_EQ:
1632+
case options::OPT_fno_profile_instr_generate:
1633+
return true;
1634+
}
1635+
return false;
1636+
}
1637+
16211638
SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
16221639
const ToolChain &HostTC, const ArgList &Args)
16231640
: ToolChain(D, Triple, Args), HostTC(HostTC),
@@ -1629,6 +1646,9 @@ SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
16291646
// Diagnose unsupported options only once.
16301647
for (OptSpecifier Opt : getUnsupportedOpts()) {
16311648
if (const Arg *A = Args.getLastArg(Opt)) {
1649+
// Native CPU can support options unsupported by other targets.
1650+
if (SupportedByNativeCPU(*this, Opt))
1651+
continue;
16321652
// All sanitizer options are not currently supported, except
16331653
// AddressSanitizer
16341654
if (A->getOption().getID() == options::OPT_fsanitize_EQ &&
@@ -1669,6 +1689,9 @@ SYCLToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
16691689
bool Unsupported = false;
16701690
for (OptSpecifier UnsupportedOpt : getUnsupportedOpts()) {
16711691
if (Opt.matches(UnsupportedOpt)) {
1692+
// NativeCPU should allow most normal cpu options.
1693+
if (SupportedByNativeCPU(*this, Opt.getID()))
1694+
continue;
16721695
if (Opt.getID() == options::OPT_fsanitize_EQ &&
16731696
A->getValues().size() == 1) {
16741697
std::string SanitizeVal = A->getValue();

clang/lib/Sema/SemaSYCL.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3625,8 +3625,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
36253625
BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
36263626
FinalizeStmts.end());
36273627

3628+
SourceLocation LL = NewBody ? NewBody->getBeginLoc() : SourceLocation();
3629+
SourceLocation LR = NewBody ? NewBody->getEndLoc() : SourceLocation();
3630+
36283631
return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts,
3629-
FPOptionsOverride(), {}, {});
3632+
FPOptionsOverride(), LL, LR);
36303633
}
36313634

36323635
void annotateHierarchicalParallelismAPICalls() {

clang/test/Driver/sycl-native-cpu.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,3 +24,9 @@
2424

2525
// RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -### 2>&1 | FileCheck -check-prefix=CHECK-NONATIVECPU %s
2626
// CHECK-NONATIVECPU-NOT: "-D" "__SYCL_NATIVE_CPU__"
27+
28+
// Checking that coverage testing options are accepted by native_cpu, and that device and host compilation invocations receive the same options
29+
// RUN: %clangxx -fsycl -fsycl-targets=native_cpu -Werror -fno-profile-instr-generate -fprofile-instr-generate -fno-coverage-mapping -fcoverage-mapping -### %s 2>&1 | FileCheck %s --check-prefix=CHECK_COV_INVO
30+
// CHECK_COV_INVO:{{.*}}clang{{.*}}-fsycl-is-device{{.*}}"-fsycl-is-native-cpu" "-D" "__SYCL_NATIVE_CPU__"{{.*}}"-fprofile-instrument=clang"{{.*}}"-fcoverage-mapping" "-fcoverage-compilation-dir={{.*}}"
31+
// CHECK_COV_INVO:{{.*}}clang{{.*}}"-fsycl-is-host"{{.*}}"-fprofile-instrument=clang"{{.*}}"-fcoverage-mapping" "-fcoverage-compilation-dir={{.*}}"
32+
Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
2+
//
3+
// Checks that the compound statement of the implicitly generated kernel body
4+
// has a valid source location (containing "line"). Previously this location
5+
// was invalid containing "<<invalid sloc>>" which causes asserts in the
6+
// llvm profiling tools.
7+
8+
#include "Inputs/sycl.hpp"
9+
10+
struct Functor {
11+
void operator()() const {}
12+
};
13+
14+
// CHECK: FunctionDecl {{.*}} _ZTS7Functor 'void ()'
15+
// CHECK-NEXT: |-CompoundStmt {{.*}} <{{.*}}line{{.*}}>
16+
17+
int main() {
18+
19+
sycl::queue().submit([&](sycl::handler &cgh) {
20+
cgh.single_task(Functor{});
21+
});
22+
}

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-7dad678",
5-
"version": "7dad678",
6-
"updated_at": "2024-11-24T10:48:51Z",
7-
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2229466354/zip",
4+
"github_tag": "igc-dev-6ee988a",
5+
"version": "6ee988a",
6+
"updated_at": "2024-11-26T15:44:10Z",
7+
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2239640503/zip",
88
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
99
}
1010
}

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,10 @@ constexpr char GENX_KERNEL_METADATA[] = "genx.kernels";
2525
// sycl/ext/oneapi/experimental/invoke_simd.hpp::__builtin_invoke_simd
2626
// overloads instantiations:
2727
constexpr char INVOKE_SIMD_PREF[] = "_Z33__regcall3____builtin_invoke_simd";
28+
// The regexp for ESIMD intrinsics:
29+
// /^_Z(\d+)__esimd_\w+/
30+
static constexpr char ESIMD_INTRIN_PREF0[] = "_Z";
31+
static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_";
2832

2933
bool isSlmAllocatorConstructor(const Function &F);
3034
bool isSlmAllocatorDestructor(const Function &F);
@@ -133,5 +137,9 @@ struct UpdateUint64MetaDataToMaxValue {
133137
// functions has changed its attribute to alwaysinline.
134138
bool prepareForAlwaysInliner(Module &M);
135139

140+
// Remove mangling from an ESIMD intrinsic function.
141+
// Returns empty on pattern match failure.
142+
StringRef stripMangling(StringRef FName);
143+
136144
} // namespace esimd
137145
} // namespace llvm

llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -129,6 +129,16 @@ void UpdateUint64MetaDataToMaxValue::operator()(Function *F) const {
129129
Node->replaceOperandWith(Key, getMetadata(New));
130130
}
131131
}
132+
StringRef stripMangling(StringRef FName) {
133+
134+
// See if the Name represents an ESIMD intrinsic and demangle only if it
135+
// does.
136+
if (!FName.consume_front(ESIMD_INTRIN_PREF0))
137+
return "";
138+
// now skip the digits
139+
FName = FName.drop_while([](char C) { return std::isdigit(C); });
140+
return FName.starts_with("__esimd") ? FName : "";
141+
}
132142

133143
} // namespace esimd
134144
} // namespace llvm

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

Lines changed: 4 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -130,10 +130,6 @@ enum class lsc_subopcode : uint8_t {
130130
read_state_info = 0x1e,
131131
fence = 0x1f,
132132
};
133-
// The regexp for ESIMD intrinsics:
134-
// /^_Z(\d+)__esimd_\w+/
135-
static constexpr char ESIMD_INTRIN_PREF0[] = "_Z";
136-
static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_";
137133
static constexpr char ESIMD_INSERTED_VSTORE_FUNC_NAME[] = "_Z14__esimd_vstorev";
138134
static constexpr char SPIRV_INTRIN_PREF[] = "__spirv_BuiltIn";
139135
struct ESIMDIntrinDesc {
@@ -2178,12 +2174,11 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
21782174
}
21792175
StringRef Name = Callee->getName();
21802176

2181-
// See if the Name represents an ESIMD intrinsic and demangle only if it
2182-
// does.
2183-
if (!Name.consume_front(ESIMD_INTRIN_PREF0) && !isDevicelibFunction(Name))
2177+
if (!isDevicelibFunction(Name))
2178+
Name = stripMangling(Name);
2179+
2180+
if (Name.empty())
21842181
continue;
2185-
// now skip the digits
2186-
Name = Name.drop_while([](char C) { return std::isdigit(C); });
21872182

21882183
// process ESIMD builtins that go through special handling instead of
21892184
// the translation procedure

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

Lines changed: 16 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,10 +8,11 @@
88
// Finds and adds sycl_explicit_simd attributes to wrapper functions that wrap
99
// ESIMD kernel functions
1010

11+
#include "llvm/IR/InstIterator.h"
12+
#include "llvm/IR/Module.h"
1113
#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
1214
#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
1315
#include "llvm/SYCLLowerIR/SYCLUtils.h"
14-
#include "llvm/IR/Module.h"
1516

1617
#define DEBUG_TYPE "LowerESIMDKernelAttrs"
1718

@@ -34,7 +35,20 @@ PreservedAnalyses
3435
SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) {
3536
bool Modified = false;
3637
for (Function &F : M) {
37-
if (llvm::esimd::isESIMD(F)) {
38+
bool ShouldConsiderESIMD = llvm::esimd::isESIMD(F);
39+
if (!ShouldConsiderESIMD) {
40+
for (Instruction &I : instructions(F)) {
41+
auto *CI = dyn_cast_or_null<CallInst>(&I);
42+
if (!CI)
43+
continue;
44+
auto *CalledF = CI->getCalledFunction();
45+
if (CalledF && !esimd::stripMangling(CalledF->getName()).empty()) {
46+
ShouldConsiderESIMD = true;
47+
break;
48+
}
49+
}
50+
}
51+
if (ShouldConsiderESIMD) {
3852
// TODO: Keep track of traversed functions to avoid repeating traversals
3953
// over same function.
4054
sycl::utils::traverseCallgraphUp(

0 commit comments

Comments
 (0)