Skip to content

Commit af9580e

Browse files
KornevNikitaKornevNikita
authored
[SYCL] Merge sycl_declared_aspects to sycl_used_aspects (#7419)
To make RT be able to throw an exception if device doesn't support the feature pass should also propagate sycl_declared aspects by merging it with sycl_used_aspects. Design: #7415
1 parent 5a188c7 commit af9580e

File tree

5 files changed

+192
-28
lines changed

5 files changed

+192
-28
lines changed

‎clang/test/CodeGenSYCL/device_has.cpp

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,32 +1,38 @@
11
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
22

3-
// Tests for IR of device_has(aspect, ...) attribute
3+
// Tests for IR of device_has(aspect, ...) attribute and
4+
// !sycl_used_aspects metadata
45
#include "sycl.hpp"
56

67
using namespace sycl;
78
queue q;
89

910
// CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
1011

11-
// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
12+
// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]]
13+
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
1214
[[sycl::device_has(sycl::aspect::cpu)]] void func1() {}
1315

14-
// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] {
16+
// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]]
17+
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS2]]
1518
[[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {}
1619

1720
// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
1821
[[sycl::device_has()]] void func3() {}
1922

20-
// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] {
23+
// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]]
24+
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS3]]
2125
template <sycl::aspect Aspect>
2226
[[sycl::device_has(Aspect)]] void func4() {}
2327

24-
// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] {
28+
// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]]
29+
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
2530
[[sycl::device_has(sycl::aspect::cpu)]] void func5();
2631
void func5() {}
2732

2833
constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
29-
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] {
34+
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]]
35+
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
3036
[[sycl::device_has(getAspect())]] void func6() {}
3137

3238
class KernelFunctor {

‎llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

Lines changed: 37 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -407,39 +407,43 @@ void propagateAspectsThroughCG(Function *F, CallGraphTy &CG,
407407
/// - checks if return and argument types are using any aspects
408408
/// - checks if instructions are using any aspects
409409
/// - updates call graph information
410-
/// - checks if function has "!sycl_used_aspects" metadata
411-
///
412-
void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects,
410+
/// - checks if function has "!sycl_used_aspects" and "!sycl_declared_aspects"
411+
/// metadata and if so collects aspects from this metadata
412+
void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToUsedAspects,
413+
FunctionToAspectsMapTy &FunctionToDeclaredAspects,
413414
TypeToAspectsMapTy &TypesWithAspects, CallGraphTy &CG) {
414415
const AspectsSetTy RetTyAspects =
415416
getAspectsFromType(F.getReturnType(), TypesWithAspects);
416-
FunctionToAspects[&F].insert(RetTyAspects.begin(), RetTyAspects.end());
417+
FunctionToUsedAspects[&F].insert(RetTyAspects.begin(), RetTyAspects.end());
417418
for (Argument &Arg : F.args()) {
418419
const AspectsSetTy ArgAspects =
419420
getAspectsFromType(Arg.getType(), TypesWithAspects);
420-
FunctionToAspects[&F].insert(ArgAspects.begin(), ArgAspects.end());
421+
FunctionToUsedAspects[&F].insert(ArgAspects.begin(), ArgAspects.end());
421422
}
422423

423424
for (Instruction &I : instructions(F)) {
424425
const AspectsSetTy Aspects =
425426
getAspectsUsedByInstruction(I, TypesWithAspects);
426-
FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end());
427+
FunctionToUsedAspects[&F].insert(Aspects.begin(), Aspects.end());
427428

428429
if (const auto *CI = dyn_cast<CallInst>(&I)) {
429430
if (!CI->isIndirectCall() && CI->getCalledFunction())
430431
CG[&F].insert(CI->getCalledFunction());
431432
}
432433
}
433434

434-
if (F.hasMetadata("sycl_used_aspects")) {
435-
const MDNode *MD = F.getMetadata("sycl_used_aspects");
436-
AspectsSetTy Aspects;
437-
for (const MDOperand &Op : MD->operands()) {
438-
Constant *C = cast<ConstantAsMetadata>(Op.get())->getValue();
439-
Aspects.insert(cast<ConstantInt>(C)->getSExtValue());
435+
auto CollectAspectsFromMD = [&F](const char* MDName, FunctionToAspectsMapTy &Map) {
436+
if (const MDNode *MD = F.getMetadata(MDName)) {
437+
AspectsSetTy Aspects;
438+
for (const MDOperand &Op : MD->operands()) {
439+
Constant *C = cast<ConstantAsMetadata>(Op.get())->getValue();
440+
Aspects.insert(cast<ConstantInt>(C)->getSExtValue());
441+
}
442+
Map[&F].insert(Aspects.begin(), Aspects.end());
440443
}
441-
FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end());
442-
}
444+
};
445+
CollectAspectsFromMD("sycl_used_aspects", FunctionToUsedAspects);
446+
CollectAspectsFromMD("sycl_declared_aspects", FunctionToDeclaredAspects);
443447
}
444448

445449
// Return true if the function is a SPIRV or SYCL builtin, e.g.
@@ -503,23 +507,34 @@ FunctionToAspectsMapTy
503507
buildFunctionsToAspectsMap(Module &M, TypeToAspectsMapTy &TypesWithAspects,
504508
const AspectValueToNameMapTy &AspectValues,
505509
const std::vector<Function *> &EntryPoints) {
506-
FunctionToAspectsMapTy FunctionToAspects;
510+
FunctionToAspectsMapTy FunctionToUsedAspects;
511+
FunctionToAspectsMapTy FunctionToDeclaredAspects;
507512
CallGraphTy CG;
508513

509514
for (Function &F : M.functions()) {
510515
if (F.isDeclaration())
511516
continue;
512-
processFunction(F, FunctionToAspects, TypesWithAspects, CG);
517+
processFunction(F, FunctionToUsedAspects, FunctionToDeclaredAspects,
518+
TypesWithAspects, CG);
513519
}
514520

515521
SmallPtrSet<const Function *, 16> Visited;
516522
for (Function *F : EntryPoints)
517-
propagateAspectsThroughCG(F, CG, FunctionToAspects, Visited);
523+
propagateAspectsThroughCG(F, CG, FunctionToUsedAspects, Visited);
524+
525+
validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues,
526+
EntryPoints, CG);
518527

519-
validateUsedAspectsForFunctions(FunctionToAspects, AspectValues, EntryPoints,
520-
CG);
528+
// The set of aspects from FunctionToDeclaredAspects should be merged to the
529+
// set of FunctionToUsedAspects after validateUsedAspectsForFunctions call to
530+
// avoid errors during validation.
531+
Visited.clear();
532+
for (Function *F : EntryPoints)
533+
propagateAspectsThroughCG(F, CG, FunctionToDeclaredAspects, Visited);
534+
for (const auto &It : FunctionToDeclaredAspects)
535+
FunctionToUsedAspects[It.first].insert(It.second.begin(), It.second.end());
521536

522-
return FunctionToAspects;
537+
return FunctionToUsedAspects;
523538
}
524539

525540
} // anonymous namespace
@@ -550,10 +565,10 @@ SYCLPropagateAspectsUsagePass::run(Module &M, ModuleAnalysisManager &MAM) {
550565

551566
propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues);
552567

553-
FunctionToAspectsMapTy FunctionToAspects = buildFunctionsToAspectsMap(
568+
FunctionToAspectsMapTy FunctionToUsedAspects = buildFunctionsToAspectsMap(
554569
M, TypesWithAspects, AspectValues, EntryPoints);
555570

556-
createUsedAspectsMetadataForFunctions(FunctionToAspects);
571+
createUsedAspectsMetadataForFunctions(FunctionToUsedAspects);
557572

558573
setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues);
559574

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-1.ll

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s
2+
3+
; kernel()
4+
; |
5+
; v
6+
; baz()
7+
; |
8+
; v
9+
; bar()
10+
; |
11+
; v
12+
; foo()
13+
14+
source_filename = "main.cpp"
15+
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"
16+
target triple = "spir64-unknown-unknown"
17+
18+
; CHECK: void @kernel() !sycl_used_aspects ![[#ASPECT:]]
19+
define weak_odr dso_local spir_kernel void @kernel() {
20+
entry:
21+
call spir_func void @_Z3bazv()
22+
ret void
23+
}
24+
25+
; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT]] {
26+
define dso_local spir_func void @_Z3bazv() {
27+
entry:
28+
call spir_func void @_Z3barv()
29+
ret void
30+
}
31+
32+
; CHECK: void @_Z3barv() !sycl_used_aspects ![[#ASPECT]] {
33+
define dso_local spir_func void @_Z3barv() {
34+
entry:
35+
call spir_func void @_Z3foov()
36+
ret void
37+
}
38+
39+
; CHECK: void @_Z3foov() !sycl_declared_aspects ![[#ASPECT]] !sycl_used_aspects ![[#ASPECT]] {
40+
define dso_local spir_func void @_Z3foov() !sycl_declared_aspects !2 {
41+
entry:
42+
ret void
43+
}
44+
45+
!sycl_aspects = !{!0, !1}
46+
47+
!0 = !{!"gpu", i32 2}
48+
!1 = !{!"fp64", i32 6}
49+
!2 = !{i32 2}

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-2.ll

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s
2+
3+
; baz()
4+
; / \
5+
; v v
6+
; bar() foo()
7+
8+
source_filename = "main.cpp"
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT1:]]
13+
define dso_local spir_kernel void @_Z3bazv() {
14+
entry:
15+
call spir_func void @_Z3barv()
16+
call spir_func void @_Z3foov()
17+
ret void
18+
}
19+
20+
; CHECK: void @_Z3barv() !sycl_used_aspects ![[#ASPECT2:]] {
21+
define dso_local spir_func void @_Z3barv() !sycl_used_aspects !3 {
22+
entry:
23+
ret void
24+
}
25+
26+
; CHECK: void @_Z3foov() !sycl_used_aspects ![[#ASPECT3:]]
27+
; CHECK-SAME: !sycl_declared_aspects ![[#ASPECT3]] {
28+
define dso_local spir_func void @_Z3foov() !sycl_declared_aspects !4 {
29+
entry:
30+
ret void
31+
}
32+
33+
; CHECK: ![[#ASPECT1]] = !{i32 2, i32 1}
34+
; CHECK: ![[#ASPECT2]] = !{i32 2}
35+
; CHECK: ![[#ASPECT3]] = !{i32 1}
36+
37+
!sycl_aspects = !{!0, !1, !2}
38+
39+
!0 = !{!"cpu", i32 1}
40+
!1 = !{!"gpu", i32 2}
41+
!2 = !{!"fp64", i32 6}
42+
!3 = !{i32 2}
43+
!4 = !{i32 1}

‎llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-3.ll

Lines changed: 51 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,51 @@
1+
; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s
2+
3+
; K
4+
; / \
5+
; F1 F2
6+
; \ / \
7+
; F3 F4
8+
9+
; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]]
10+
define spir_kernel void @kernel() {
11+
call spir_func void @func1()
12+
call spir_func void @func2()
13+
ret void
14+
}
15+
16+
; CHECK: spir_func void @func1() !sycl_used_aspects ![[#ID2:]] {
17+
define spir_func void @func1() {
18+
call spir_func void @func3()
19+
ret void
20+
}
21+
22+
; CHECK: spir_func void @func2() !sycl_used_aspects ![[#ID1]] {
23+
define spir_func void @func2() {
24+
call spir_func void @func3()
25+
call spir_func void @func4()
26+
ret void
27+
}
28+
29+
; CHECK: spir_func void @func3() !sycl_used_aspects ![[#ID2]] {
30+
define spir_func void @func3() !sycl_used_aspects !4 {
31+
ret void
32+
}
33+
34+
; CHECK: spir_func void @func4() !sycl_used_aspects ![[#ID3:]]
35+
; CHECK-SAME: !sycl_declared_aspects ![[#ID3]] {
36+
define spir_func void @func4() !sycl_declared_aspects !3 {
37+
ret void
38+
}
39+
40+
!sycl_aspects = !{!0, !1, !2}
41+
42+
!0 = !{!"host", i32 0}
43+
!1 = !{!"cpu", i32 1}
44+
!2 = !{!"fp64", i32 6}
45+
!3 = !{i32 0}
46+
!4 = !{i32 1}
47+
!5 = !{i32 0, i32 1}
48+
49+
; CHECK: ![[#ID1]] = !{i32 1, i32 0}
50+
; CHECK: ![[#ID2]] = !{i32 1}
51+
; CHECK: ![[#ID3]] = !{i32 0}

0 commit comments

Comments
 (0)