Skip to content

Commit

Permalink
[SYCL] Merge sycl_declared_aspects to sycl_used_aspects (#7419)
Browse files Browse the repository at this point in the history 
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
  • Loading branch information
@KornevNikita
KornevNikita authored Nov 17, 2022
1 parent 5a188c7 commit af9580e
Show file tree
Hide file tree
Showing 5 changed files with 192 additions and 28 deletions.
18 changes: 12 additions & 6 deletions clang/test/CodeGenSYCL/device_has.cpp
View file
Original file line number Diff line number Diff line change
@@ -1,32 +1,38 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

// Tests for IR of device_has(aspect, ...) attribute
// Tests for IR of device_has(aspect, ...) attribute and
// !sycl_used_aspects metadata
#include "sycl.hpp"

using namespace sycl;
queue q;

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

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

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

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

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

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

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

class KernelFunctor {
Expand Down
59 changes: 37 additions & 22 deletions llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -407,39 +407,43 @@ void propagateAspectsThroughCG(Function *F, CallGraphTy &CG,
/// - checks if return and argument types are using any aspects
/// - checks if instructions are using any aspects
/// - updates call graph information
/// - checks if function has "!sycl_used_aspects" metadata
///
void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToAspects,
/// - checks if function has "!sycl_used_aspects" and "!sycl_declared_aspects"
/// metadata and if so collects aspects from this metadata
void processFunction(Function &F, FunctionToAspectsMapTy &FunctionToUsedAspects,
FunctionToAspectsMapTy &FunctionToDeclaredAspects,
TypeToAspectsMapTy &TypesWithAspects, CallGraphTy &CG) {
const AspectsSetTy RetTyAspects =
getAspectsFromType(F.getReturnType(), TypesWithAspects);
FunctionToAspects[&F].insert(RetTyAspects.begin(), RetTyAspects.end());
FunctionToUsedAspects[&F].insert(RetTyAspects.begin(), RetTyAspects.end());
for (Argument &Arg : F.args()) {
const AspectsSetTy ArgAspects =
getAspectsFromType(Arg.getType(), TypesWithAspects);
FunctionToAspects[&F].insert(ArgAspects.begin(), ArgAspects.end());
FunctionToUsedAspects[&F].insert(ArgAspects.begin(), ArgAspects.end());
}

for (Instruction &I : instructions(F)) {
const AspectsSetTy Aspects =
getAspectsUsedByInstruction(I, TypesWithAspects);
FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end());
FunctionToUsedAspects[&F].insert(Aspects.begin(), Aspects.end());

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

if (F.hasMetadata("sycl_used_aspects")) {
const MDNode *MD = F.getMetadata("sycl_used_aspects");
AspectsSetTy Aspects;
for (const MDOperand &Op : MD->operands()) {
Constant *C = cast<ConstantAsMetadata>(Op.get())->getValue();
Aspects.insert(cast<ConstantInt>(C)->getSExtValue());
auto CollectAspectsFromMD = [&F](const char* MDName, FunctionToAspectsMapTy &Map) {
if (const MDNode *MD = F.getMetadata(MDName)) {
AspectsSetTy Aspects;
for (const MDOperand &Op : MD->operands()) {
Constant *C = cast<ConstantAsMetadata>(Op.get())->getValue();
Aspects.insert(cast<ConstantInt>(C)->getSExtValue());
}
Map[&F].insert(Aspects.begin(), Aspects.end());
}
FunctionToAspects[&F].insert(Aspects.begin(), Aspects.end());
}
};
CollectAspectsFromMD("sycl_used_aspects", FunctionToUsedAspects);
CollectAspectsFromMD("sycl_declared_aspects", FunctionToDeclaredAspects);
}

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

for (Function &F : M.functions()) {
if (F.isDeclaration())
continue;
processFunction(F, FunctionToAspects, TypesWithAspects, CG);
processFunction(F, FunctionToUsedAspects, FunctionToDeclaredAspects,
TypesWithAspects, CG);
}

SmallPtrSet<const Function *, 16> Visited;
for (Function *F : EntryPoints)
propagateAspectsThroughCG(F, CG, FunctionToAspects, Visited);
propagateAspectsThroughCG(F, CG, FunctionToUsedAspects, Visited);

validateUsedAspectsForFunctions(FunctionToUsedAspects, AspectValues,
EntryPoints, CG);

validateUsedAspectsForFunctions(FunctionToAspects, AspectValues, EntryPoints,
CG);
// The set of aspects from FunctionToDeclaredAspects should be merged to the
// set of FunctionToUsedAspects after validateUsedAspectsForFunctions call to
// avoid errors during validation.
Visited.clear();
for (Function *F : EntryPoints)
propagateAspectsThroughCG(F, CG, FunctionToDeclaredAspects, Visited);
for (const auto &It : FunctionToDeclaredAspects)
FunctionToUsedAspects[It.first].insert(It.second.begin(), It.second.end());

return FunctionToAspects;
return FunctionToUsedAspects;
}

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

propagateAspectsToOtherTypesInModule(M, TypesWithAspects, AspectValues);

FunctionToAspectsMapTy FunctionToAspects = buildFunctionsToAspectsMap(
FunctionToAspectsMapTy FunctionToUsedAspects = buildFunctionsToAspectsMap(
M, TypesWithAspects, AspectValues, EntryPoints);

createUsedAspectsMetadataForFunctions(FunctionToAspects);
createUsedAspectsMetadataForFunctions(FunctionToUsedAspects);

setSyclFixedTargetsMD(EntryPoints, TargetFixedAspects, AspectValues);

Expand Down
49 changes: 49 additions & 0 deletions llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-1.ll
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s

; kernel()
; |
; v
; baz()
; |
; v
; bar()
; |
; v
; foo()

source_filename = "main.cpp"
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"
target triple = "spir64-unknown-unknown"

; CHECK: void @kernel() !sycl_used_aspects ![[#ASPECT:]]
define weak_odr dso_local spir_kernel void @kernel() {
entry:
call spir_func void @_Z3bazv()
ret void
}

; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT]] {
define dso_local spir_func void @_Z3bazv() {
entry:
call spir_func void @_Z3barv()
ret void
}

; CHECK: void @_Z3barv() !sycl_used_aspects ![[#ASPECT]] {
define dso_local spir_func void @_Z3barv() {
entry:
call spir_func void @_Z3foov()
ret void
}

; CHECK: void @_Z3foov() !sycl_declared_aspects ![[#ASPECT]] !sycl_used_aspects ![[#ASPECT]] {
define dso_local spir_func void @_Z3foov() !sycl_declared_aspects !2 {
entry:
ret void
}

!sycl_aspects = !{!0, !1}

!0 = !{!"gpu", i32 2}
!1 = !{!"fp64", i32 6}
!2 = !{i32 2}
43 changes: 43 additions & 0 deletions llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-2.ll
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,43 @@
; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s

; baz()
; / \
; v v
; bar() foo()

source_filename = "main.cpp"
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"
target triple = "spir64-unknown-unknown"

; CHECK: void @_Z3bazv() !sycl_used_aspects ![[#ASPECT1:]]
define dso_local spir_kernel void @_Z3bazv() {
entry:
call spir_func void @_Z3barv()
call spir_func void @_Z3foov()
ret void
}

; CHECK: void @_Z3barv() !sycl_used_aspects ![[#ASPECT2:]] {
define dso_local spir_func void @_Z3barv() !sycl_used_aspects !3 {
entry:
ret void
}

; CHECK: void @_Z3foov() !sycl_used_aspects ![[#ASPECT3:]]
; CHECK-SAME: !sycl_declared_aspects ![[#ASPECT3]] {
define dso_local spir_func void @_Z3foov() !sycl_declared_aspects !4 {
entry:
ret void
}

; CHECK: ![[#ASPECT1]] = !{i32 2, i32 1}
; CHECK: ![[#ASPECT2]] = !{i32 2}
; CHECK: ![[#ASPECT3]] = !{i32 1}

!sycl_aspects = !{!0, !1, !2}

!0 = !{!"cpu", i32 1}
!1 = !{!"gpu", i32 2}
!2 = !{!"fp64", i32 6}
!3 = !{i32 2}
!4 = !{i32 1}
51 changes: 51 additions & 0 deletions llvm/test/SYCLLowerIR/PropagateAspectsUsage/PropageteDeclaredAspects/propagate-declared-3.ll
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,51 @@
; RUN: opt -passes=sycl-propagate-aspects-usage %s -S | FileCheck %s

; K
; / \
; F1 F2
; \ / \
; F3 F4

; CHECK: spir_kernel void @kernel() !sycl_used_aspects ![[#ID1:]]
define spir_kernel void @kernel() {
call spir_func void @func1()
call spir_func void @func2()
ret void
}

; CHECK: spir_func void @func1() !sycl_used_aspects ![[#ID2:]] {
define spir_func void @func1() {
call spir_func void @func3()
ret void
}

; CHECK: spir_func void @func2() !sycl_used_aspects ![[#ID1]] {
define spir_func void @func2() {
call spir_func void @func3()
call spir_func void @func4()
ret void
}

; CHECK: spir_func void @func3() !sycl_used_aspects ![[#ID2]] {
define spir_func void @func3() !sycl_used_aspects !4 {
ret void
}

; CHECK: spir_func void @func4() !sycl_used_aspects ![[#ID3:]]
; CHECK-SAME: !sycl_declared_aspects ![[#ID3]] {
define spir_func void @func4() !sycl_declared_aspects !3 {
ret void
}

!sycl_aspects = !{!0, !1, !2}

!0 = !{!"host", i32 0}
!1 = !{!"cpu", i32 1}
!2 = !{!"fp64", i32 6}
!3 = !{i32 0}
!4 = !{i32 1}
!5 = !{i32 0, i32 1}

; CHECK: ![[#ID1]] = !{i32 1, i32 0}
; CHECK: ![[#ID2]] = !{i32 1}
; CHECK: ![[#ID3]] = !{i32 0}

0 comments on commit af9580e

Please sign in to comment.