Skip to content

Commit 81882ed

Browse files
KornevNikitaAlexeySachkov
KornevNikita
and
AlexeySachkov
authored
[SYCL] Propagate sycl_fixed_targets metadata (#7115)
As part of optional kernel features implementation we need to implement new compiler switch option "-fsycl-fixed-targets=target,target..." to diagnose optional feature usage and mark SYCL kernel functions with a "!sycl_fixed_targets !N" metadata, where "!N" is a list of targets as aspects enum members. This patch applies changes for SYCLPropagateAspectsUsage pass and prepare it to accept new switch from Clang. Co-authored-by: Alexey Sachkov <[email protected]>
1 parent ea7e716 commit 81882ed

File tree

12 files changed

+89
-22
lines changed

12 files changed

+89
-22
lines changed

‎llvm/include/llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,14 @@ namespace llvm {
2020
class SYCLPropagateAspectsUsagePass
2121
: public PassInfoMixin<SYCLPropagateAspectsUsagePass> {
2222
public:
23+
SYCLPropagateAspectsUsagePass(StringRef OptionsString = {}) {
24+
OptionsString.split(this->TargetFixedAspects, ',', /*MaxSplit=*/-1,
25+
/*KeepEmpty=*/false);
26+
};
2327
PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
28+
29+
private:
30+
SmallVector<StringRef, 8> TargetFixedAspects;
2431
};
2532

2633
} // namespace llvm

‎llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp‎

Lines changed: 47 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,7 @@
3737
#include "llvm/IR/IntrinsicInst.h"
3838
#include "llvm/IR/Module.h"
3939
#include "llvm/Pass.h"
40+
#include "llvm/Support/CommandLine.h"
4041
#include "llvm/Support/Path.h"
4142

4243
#include <queue>
@@ -45,6 +46,12 @@
4546

4647
using namespace llvm;
4748

49+
static cl::opt<std::string> ClSyclFixedTargets(
50+
"sycl-propagate-aspects-usage-fixed-targets",
51+
cl::desc("Specify target device(s) all device code in the translation unit "
52+
"is expected to be runnable on"),
53+
cl::Hidden, cl::init(""));
54+
4855
namespace {
4956

5057
using AspectsSetTy = SmallSet<int, 4>;
@@ -348,19 +355,41 @@ bool isEntryPoint(const Function &F) {
348355
return F.hasFnAttribute("sycl-module-id") && !isSpirvSyclBuiltin(F.getName());
349356
}
350357

358+
void setSyclFixedTargetsMD(const std::vector<Function *> &EntryPoints,
359+
const SmallVector<StringRef, 8> &Targets,
360+
AspectValueToNameMapTy &AspectValues) {
361+
if (EntryPoints.empty())
362+
return;
363+
364+
SmallVector<Metadata *, 8> TargetsMD;
365+
LLVMContext &C = EntryPoints[0]->getContext();
366+
367+
for (const auto &Target : Targets) {
368+
if (!Target.empty()) {
369+
auto AspectIt = AspectValues.find(Target);
370+
if (AspectIt != AspectValues.end()) {
371+
auto ConstIntTarget =
372+
ConstantInt::getSigned(Type::getInt32Ty(C), AspectIt->second);
373+
TargetsMD.push_back(ConstantAsMetadata::get(ConstIntTarget));
374+
}
375+
}
376+
}
377+
378+
MDNode *MDN = MDNode::get(C, TargetsMD);
379+
for (Function *F : EntryPoints)
380+
F->setMetadata("sycl_fixed_targets", MDN);
381+
}
382+
351383
/// Returns a map of functions with corresponding used aspects.
352384
FunctionToAspectsMapTy
353-
buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects) {
385+
buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
386+
const std::vector<Function *> &EntryPoints) {
354387
FunctionToAspectsMapTy FunctionToAspects;
355388
CallGraphTy CG;
356-
std::vector<Function *> EntryPoints;
389+
357390
for (Function &F : M.functions()) {
358391
if (F.isDeclaration())
359392
continue;
360-
361-
if (isEntryPoint(F))
362-
EntryPoints.push_back(&F);
363-
364393
processFunction(F, FunctionToAspects, TypesWithAspects, CG);
365394
}
366395

@@ -388,14 +417,25 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) {
388417
return PreservedAnalyses::all();
389418
}
390419

420+
if (ClSyclFixedTargets.getNumOccurrences() > 0)
421+
StringRef(ClSyclFixedTargets)
422+
.split(TargetFixedAspects, ',', /*MaxSplit=*/-1, /*KeepEmpty=*/false);
423+
424+
std::vector<Function *> EntryPoints;
425+
for (Function &F : M.functions())
426+
if (isEntryPoint(F))
427+
EntryPoints.push_back(&F);
428+
391429
propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues);
392430

393431
FunctionToAspectsMapTy FunctionToAspects =
394-
buildFunctionsToAspectsMap(M, TypesWithAspects);
432+
buildFunctionsToAspectsMap(M, TypesWithAspects, EntryPoints);
395433

396434
createUsedAspectsMetadataForFunctions(FunctionToAspects);
397435
// FIXME: check and diagnose if a function uses an aspect which was not
398436
// declared through [[sycl::device_has()]] attribute
399437

438+
setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues);
439+
400440
return PreservedAnalyses::all();
401441
}

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropagateFunctionWithSyclDetailAspect/propagate-aspects-from-function-with-sycl-detail-1.ll‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,8 +14,8 @@ source_filename = "main.cpp"
1414
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
1515
target triple = "spir64-unknown-unknown"
1616

17-
; CHECK: void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_() !sycl_used_aspects ![[#ASPECT:]] {
18-
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_() {
17+
; CHECK: void @kernel() !sycl_used_aspects ![[#ASPECT:]]
18+
define weak_odr dso_local spir_kernel void @kernel() {
1919
entry:
2020
call spir_func void @_Z3bazv()
2121
ret void

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropagateFunctionWithSyclDetailAspect/propagate-aspects-from-function-with-sycl-detail-2.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@ source_filename = "main.cpp"
1111
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
1212
target triple = "spir64-unknown-unknown"
1313

14-
; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT1:]] {
14+
; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT1:]]
1515
define dso_local spir_kernel void @_Z3bazv() {
1616
entry:
1717
call spir_func void @_Z3barv()

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropagateFunctionWithSyclDetailAspect/propagate-aspects-from-function-with-sycl-detail-3.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@
1515
%Optional.A = type { i32 }
1616
%Optional.B = type { i32 }
1717

18-
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]] {
18+
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]]
1919
define spir_kernel void @kernel() {
2020
call spir_func void @func1()
2121
call spir_func void @func2()

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-1.ll‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,14 +14,14 @@
1414
%Optional.A = type { i32 }
1515
%Optional.B = type { i32 }
1616

17-
; CHECK: spir_kernel void @kernel1() !sycl_used_aspects ![[#ID1:]] {
17+
; CHECK: spir_kernel void @kernel1() !sycl_used_aspects ![[#ID1:]]
1818
define spir_kernel void @kernel1() {
1919
call spir_func void @func1()
2020
call spir_func void @func2()
2121
ret void
2222
}
2323

24-
; CHECK: spir_kernel void @kernel2() !sycl_used_aspects ![[#ID2:]] {
24+
; CHECK: spir_kernel void @kernel2() !sycl_used_aspects ![[#ID2:]]
2525
define spir_kernel void @kernel2() {
2626
call spir_func void @func2()
2727
call spir_func void @func3()

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/call-graph-2.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@
1515
%Optional.A = type { i32 }
1616
%Optional.B = type { i32 }
1717

18-
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]] {
18+
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]]
1919
define spir_kernel void @kernel() {
2020
call spir_func void @func1()
2121
call spir_func void @func2()

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/composite-types-1.ll‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@
2222

2323
%F2.does.not.contain.optional = type { %B.core, %C.core*, %D2.does.not.contain.optional* }
2424

25-
; CHECK: spir_kernel void @kernelD1.uses.optional() !sycl_used_aspects ![[MDID:[0-9]+]] {
25+
; CHECK: spir_kernel void @kernelD1.uses.optional() !sycl_used_aspects ![[MDID:[0-9]+]]
2626
define spir_kernel void @kernelD1.uses.optional() {
2727
%tmp = alloca %D1.contains.optional
2828
ret void
@@ -34,7 +34,7 @@ define spir_func void @funcD1.uses.optional() {
3434
ret void
3535
}
3636

37-
; CHECK: spir_kernel void @kernelD2.does.not.use.optional() {
37+
; CHECK: spir_kernel void @kernelD2.does.not.use.optional()
3838
define spir_kernel void @kernelD2.does.not.use.optional() {
3939
%tmp = alloca %D2.does.not.contain.optional
4040
ret void
@@ -46,7 +46,7 @@ define spir_func void @funcD2.does.not.use.optional() {
4646
ret void
4747
}
4848

49-
; CHECK: spir_kernel void @kernelE.uses.optional() !sycl_used_aspects ![[MDID]] {
49+
; CHECK: spir_kernel void @kernelE.uses.optional() !sycl_used_aspects ![[MDID]]
5050
define spir_kernel void @kernelE.uses.optional() {
5151
%tmp = alloca %E.contains.optional
5252
ret void
@@ -58,7 +58,7 @@ define spir_func void @funcE.uses.optional() {
5858
ret void
5959
}
6060

61-
; CHECK: spir_kernel void @kernelF1.points.to.optional() {
61+
; CHECK: spir_kernel void @kernelF1.points.to.optional()
6262
define spir_kernel void @kernelF1.points.to.optional() {
6363
%tmp = alloca %F1.points.to.optional
6464
ret void
@@ -70,7 +70,7 @@ define spir_func void @funcF1.points.to.optional() {
7070
ret void
7171
}
7272

73-
; CHECK: spir_kernel void @kernelF2.does.not.use.optional() {
73+
; CHECK: spir_kernel void @kernelF2.does.not.use.optional()
7474
define spir_kernel void @kernelF2.does.not.use.optional() {
7575
%tmp = alloca %F2.does.not.contain.optional
7676
ret void

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/double.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44

55
%composite = type { double }
66

7-
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[MDID:[0-9]+]] {
7+
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[MDID:[0-9]+]]
88
define spir_kernel void @kernel() {
99
call spir_func void @func()
1010
ret void

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/multiple-aspects.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@ define spir_func void @funcD() {
3434
ret void
3535
}
3636

37-
; CHECK: define spir_kernel void @kernel() !sycl_used_aspects ![[#ID3]] {
37+
; CHECK: define spir_kernel void @kernel() !sycl_used_aspects ![[#ID3]]
3838
define spir_kernel void @kernel() {
3939
call spir_func void @funcD()
4040
ret void

0 commit comments

Comments
 (0)