Skip to content

[SYCL] Merge sycl_declared_aspects to sycl_used_aspects #7419

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 8 commits into from
Nov 17, 2022
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 12 additions & 6 deletions clang/test/CodeGenSYCL/device_has.cpp
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]+]]
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we run the pass even if -disable-llvm-passes option is passed?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm trying to understand why the clang-only test was changed :)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hm, I fixed this test because it was failed. Seems like we are.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Seems like we need this FE test changes for now

Copy link
Contributor

@elizabethandrews elizabethandrews Nov 16, 2022

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I assume sycl_used_aspects metadata is being added here since the pass runs even if we pass -disable-llvm-passes. Can you modify the CHECK to test for it as well?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Exactly. Sure thing

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done 435c611

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please modify the comment at the top of the test to indicate we check sycl_used_aspects here as well.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

// 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
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
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}
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}
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}