Skip to content

Commit 5d92897

Browse files
authoredJul 10, 2023
[SYCL][clang] Fix uses_aspects applied to function declarations (#10164)
It was only applied to definitions.
·
v6.1.02023-WW46
1 parent fda9171 commit 5d92897

File tree

3 files changed

+27
-23
lines changed

3 files changed

+27
-23
lines changed
 

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 0 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1121,16 +1121,6 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
11211121
Fn->setMetadata("loop_fuse",
11221122
llvm::MDNode::get(getLLVMContext(), AttrMDArgs));
11231123
}
1124-
if (const auto *A = D->getAttr<SYCLUsesAspectsAttr>()) {
1125-
SmallVector<llvm::Metadata *, 4> AspectsMD;
1126-
for (auto *Aspect : A->aspects()) {
1127-
llvm::APSInt AspectInt = Aspect->EvaluateKnownConstInt(getContext());
1128-
AspectsMD.push_back(llvm::ConstantAsMetadata::get(
1129-
Builder.getInt32(AspectInt.getZExtValue())));
1130-
}
1131-
Fn->setMetadata("sycl_used_aspects",
1132-
llvm::MDNode::get(getLLVMContext(), AspectsMD));
1133-
}
11341124

11351125
// Source location of functions is required to emit required diagnostics in
11361126
// SYCLPropagateAspectsUsagePass. Save the token in a srcloc metadata node.

‎clang/lib/CodeGen/CodeGenModule.cpp

Lines changed: 23 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -2648,6 +2648,19 @@ void CodeGenModule::finalizeKCFITypes() {
26482648
}
26492649
}
26502650

2651+
template <typename AttrT>
2652+
void applySYCLAspectsMD(AttrT *A, ASTContext &ACtx, llvm::LLVMContext &LLVMCtx,
2653+
llvm::Function *F, StringRef MDName) {
2654+
SmallVector<llvm::Metadata *, 4> AspectsMD;
2655+
for (auto *Aspect : A->aspects()) {
2656+
llvm::APSInt AspectInt = Aspect->EvaluateKnownConstInt(ACtx);
2657+
auto *T = llvm::Type::getInt32Ty(LLVMCtx);
2658+
auto *C = llvm::Constant::getIntegerValue(T, AspectInt);
2659+
AspectsMD.push_back(llvm::ConstantAsMetadata::get(C));
2660+
}
2661+
F->setMetadata(MDName, llvm::MDNode::get(LLVMCtx, AspectsMD));
2662+
}
2663+
26512664
void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
26522665
bool IsIncompleteFunction,
26532666
bool IsThunk) {
@@ -2755,6 +2768,15 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
27552768
CalleeIdx, PayloadIndices,
27562769
/* VarArgsArePassed */ false)}));
27572770
}
2771+
2772+
// Apply SYCL specific attributes/metadata.
2773+
if (const auto *A = FD->getAttr<SYCLDeviceHasAttr>())
2774+
applySYCLAspectsMD(A, getContext(), getLLVMContext(), F,
2775+
"sycl_declared_aspects");
2776+
2777+
if (const auto *A = FD->getAttr<SYCLUsesAspectsAttr>())
2778+
applySYCLAspectsMD(A, getContext(), getLLVMContext(), F,
2779+
"sycl_used_aspects");
27582780
}
27592781

27602782
void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) {
@@ -4573,20 +4595,8 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
45734595
}
45744596

45754597
assert(F->getName() == MangledName && "name was uniqued!");
4576-
if (D) {
4598+
if (D)
45774599
SetFunctionAttributes(GD, F, IsIncompleteFunction, IsThunk);
4578-
if (const auto *A = D->getAttr<SYCLDeviceHasAttr>()) {
4579-
SmallVector<llvm::Metadata *, 4> AspectsMD;
4580-
for (auto *Aspect : A->aspects()) {
4581-
llvm::APSInt AspectInt = Aspect->EvaluateKnownConstInt(getContext());
4582-
auto *T = llvm::Type::getInt32Ty(getLLVMContext());
4583-
auto *C = llvm::Constant::getIntegerValue(T, AspectInt);
4584-
AspectsMD.push_back(llvm::ConstantAsMetadata::get(C));
4585-
}
4586-
F->setMetadata("sycl_declared_aspects",
4587-
llvm::MDNode::get(getLLVMContext(), AspectsMD));
4588-
}
4589-
}
45904600
if (ExtraAttrs.hasFnAttrs()) {
45914601
llvm::AttrBuilder B(F->getContext(), ExtraAttrs.getFnAttrs());
45924602
F->addFnAttrs(B);

‎clang/test/CodeGenSYCL/uses_aspects.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,9 @@ constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
3838
// CHECK: define dso_local spir_func void @{{.*}}func7{{.*}} !sycl_used_aspects ![[ASPECTS1]] {
3939
[[__sycl_detail__::__uses_aspects__(getAspect())]] void func7() {}
4040

41+
// CHECK: declare !sycl_used_aspects ![[ASPECTS1]] spir_func void @{{.*}}func8{{.*}}
42+
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] SYCL_EXTERNAL void func8();
43+
4144
class KernelFunctor {
4245
public:
4346
void operator()() const {
@@ -48,6 +51,7 @@ class KernelFunctor {
4851
func5();
4952
func6();
5053
func7();
54+
func8();
5155
}
5256
};
5357

0 commit comments

Comments
 (0)
Please sign in to comment.