Skip to content

[SYCL][clang] Fix uses_aspects applied to function declarations #10164

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.

Sign up for GitHub

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 2 commits into from
Jul 10, 2023
Merged

[SYCL][clang] Fix uses_aspects applied to function declarations #10164

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9443d8e
[SYCL][clang] Fix uses_aspects applied to function declarations
Fznamznon Jun 30, 2023
6ce3a13
Merge branch 'intel:sycl' into private/mpodchis/fix-uses-aspects
Fznamznon Jul 6, 2023
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
10 changes: 0 additions & 10 deletions clang/lib/CodeGen/CodeGenFunction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1121,16 +1121,6 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
Fn->setMetadata("loop_fuse",
llvm::MDNode::get(getLLVMContext(), AttrMDArgs));
}
if (const auto *A = D->getAttr<SYCLUsesAspectsAttr>()) {
SmallVector<llvm::Metadata *, 4> AspectsMD;
for (auto *Aspect : A->aspects()) {
llvm::APSInt AspectInt = Aspect->EvaluateKnownConstInt(getContext());
AspectsMD.push_back(llvm::ConstantAsMetadata::get(
Builder.getInt32(AspectInt.getZExtValue())));
}
Fn->setMetadata("sycl_used_aspects",
llvm::MDNode::get(getLLVMContext(), AspectsMD));
}

// Source location of functions is required to emit required diagnostics in
// SYCLPropagateAspectsUsagePass. Save the token in a srcloc metadata node.
Expand Down
36 changes: 23 additions & 13 deletions clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2648,6 +2648,19 @@ void CodeGenModule::finalizeKCFITypes() {
}
}

template <typename AttrT>
void applySYCLAspectsMD(AttrT *A, ASTContext &ACtx, llvm::LLVMContext &LLVMCtx,
llvm::Function *F, StringRef MDName) {
SmallVector<llvm::Metadata *, 4> AspectsMD;
for (auto *Aspect : A->aspects()) {
llvm::APSInt AspectInt = Aspect->EvaluateKnownConstInt(ACtx);
auto *T = llvm::Type::getInt32Ty(LLVMCtx);
auto *C = llvm::Constant::getIntegerValue(T, AspectInt);
AspectsMD.push_back(llvm::ConstantAsMetadata::get(C));
}
F->setMetadata(MDName, llvm::MDNode::get(LLVMCtx, AspectsMD));
}

void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
bool IsIncompleteFunction,
bool IsThunk) {
Expand Down Expand Up @@ -2755,6 +2768,15 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
CalleeIdx, PayloadIndices,
/* VarArgsArePassed */ false)}));
}

// Apply SYCL specific attributes/metadata.
if (const auto *A = FD->getAttr<SYCLDeviceHasAttr>())
applySYCLAspectsMD(A, getContext(), getLLVMContext(), F,
"sycl_declared_aspects");

if (const auto *A = FD->getAttr<SYCLUsesAspectsAttr>())
applySYCLAspectsMD(A, getContext(), getLLVMContext(), F,
"sycl_used_aspects");
}

void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) {
Expand Down Expand Up @@ -4573,20 +4595,8 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
}

assert(F->getName() == MangledName && "name was uniqued!");
if (D) {
if (D)
SetFunctionAttributes(GD, F, IsIncompleteFunction, IsThunk);
if (const auto *A = D->getAttr<SYCLDeviceHasAttr>()) {
SmallVector<llvm::Metadata *, 4> AspectsMD;
for (auto *Aspect : A->aspects()) {
llvm::APSInt AspectInt = Aspect->EvaluateKnownConstInt(getContext());
auto *T = llvm::Type::getInt32Ty(getLLVMContext());
auto *C = llvm::Constant::getIntegerValue(T, AspectInt);
AspectsMD.push_back(llvm::ConstantAsMetadata::get(C));
}
F->setMetadata("sycl_declared_aspects",
llvm::MDNode::get(getLLVMContext(), AspectsMD));
}
}
if (ExtraAttrs.hasFnAttrs()) {
llvm::AttrBuilder B(F->getContext(), ExtraAttrs.getFnAttrs());
F->addFnAttrs(B);
Expand Down
4 changes: 4 additions & 0 deletions clang/test/CodeGenSYCL/uses_aspects.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,9 @@ constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
// CHECK: define dso_local spir_func void @{{.*}}func7{{.*}} !sycl_used_aspects ![[ASPECTS1]] {
[[__sycl_detail__::__uses_aspects__(getAspect())]] void func7() {}

// CHECK: declare !sycl_used_aspects ![[ASPECTS1]] spir_func void @{{.*}}func8{{.*}}
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] SYCL_EXTERNAL void func8();
Copy link
Contributor

Choose a reason for hiding this comment

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

The design document specifies this metadata is applied to function definitions - https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/design/OptionalDeviceFeatures.md

How does this work with redeclarations? What is finally attached to the function definition? IIRC there is some pass which propagates this metadata up the static call graph. What gets propagated now if we have different attribute arguments in different declarations

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The design document specifies this metadata is applied to function definitions - https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/design/OptionalDeviceFeatures.md

The document seems to be assuming that the atribute can be applied to function declarations, see for example https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/design/OptionalDeviceFeatures.md#changes-to-dpc-headers . However it says,

If a function is decorated with [[sycl_detail::uses_aspects()]], the front-end adds an !sycl_used_aspects metadata to the function's definition listing the aspects from that attribute.

If a function is decorated with [[sycl::device_has()]], the front-end adds an !sycl_declared_aspects metadata to the function's definition listing the aspects from that attribute.

For both uses_aspects and device_has attibutes. We already has support for device_has on declarations added by https://github.com/intel/llvm/pull/9611/files . I'm not sure why it is limited to definitions in the document, maybe we should change the wording to "function signature"? CC @gmlueck, @AlexeySachkov to clarify.

How does this work with redeclarations? IIRC there is some pass which propagates this metadata up the static call graph. What gets propagated now if we have different attribute arguments in different declarations

When redeclarations appear, the most recent is referenced by the code. Meaning the metadata is generated based on attirbute value applied to the most recent declarations. Even in case like

SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp64)]] void foo() ;

kernel1 calling foo

SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16)]] void foo() ;

kernel2 calling foo

I'm seeing only fp16 aspect on kernels and function declaration.

The pass seem to be updated by https://github.com/intel/llvm/pull/9611/files in generic way, so it seems to be working fine with declarations.
The metadata generated for the most recent redeclaration is propagated. When the most recent doesn't have the attribute, metadata is propagated from previous redeclaration.

The thing I find confusing for the end user, when attribute is applied to a redeclaration and the previous declaration had an attribute, the warning is emitted:

t.cpp:29:17: warning: attribute '__uses_aspects__' is already applied [-Wignored-attributes]
   29 | SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16)]] void foo();
      |                 ^
t.cpp:4:17: note: previous attribute is here
    4 | SYCL_EXTERNAL [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp64)]] void foo()
      |                 ^
1 warning generated.

It seems to be saying that the attribute is ignored on the most recent declaration, however instead it seems to be ignored on the first declaration since the metadata is generated based on the latter declaration.

Copy link
Contributor

Choose a reason for hiding this comment

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

The [[sycl::device_has]] attribute is defined in the SYCL 2020 specification section 5.8 "Attributes for device code", which says this about attributes and redeclarations:

If one of the attributes defined in this section is applied to a kernel or device function, it must be applied
to the first declaration of that kernel or device function in the translation unit. Programs which fail to do
this are ill formed and the compiler must issue a diagnostic. Redeclarations of the kernel or device function in the same translation unit may optionally have the same attribute applied (so long as the attribute
arguments are the same between the declarations), but this is not required. The attribute remains in
effect regardless of whether it appears in the redeclaration.

Based on that wording, I think it is an error to redeclare a function with a different [[sycl::device_has]] attribute, and the compiler is supposed to diagnose this case.

I think it is reasonable to treat the [sycl_detail::uses_aspects]] attribute in the same way.

BTW, the links you have above are pointing to a different company's fork of our repo. When looking at our design documents, you should look at our own fork (not the triSYCL fork). For example, this is the link to our copy of the optional device features design document:

https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md

Copy link
Contributor Author

Choose a reason for hiding this comment

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

BTW, the links you have above are pointing to a different company's fork of our repo. When looking at our design documents, you should look at our own fork (not the triSYCL fork)

OMG, I didn't notice that the first link posted by Elizabeth is to triSYCL repo. Thanks for pointing this out.

Based on that wording, I think it is an error to redeclare a function with a different [[sycl::device_has]] attribute, and the compiler is supposed to diagnose this case.

Well, apparently the compiler's behavior is different now. I would prefer implementing the diagnostic separately though.

Copy link
Contributor

Choose a reason for hiding this comment

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

BTW, the links you have above are pointing to a different company's fork of our repo.

I didn't notice I checked the wrong link! I just opened the first search item on google. I apologize!

Copy link
Contributor

Choose a reason for hiding this comment

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

It seems to be saying that the attribute is ignored on the most recent declaration, however instead it seems to be ignored on the first declaration since the metadata is generated based on the latter declaration.

I believe this is 'normal' clang attribute functionality. Warnings are usually generated on redeclarations but metadata is also usually from first declaration. So there is some bug there. Anyway based on the spec @gmlueck pasted above, it looks like we need to be parsing attribute arguments and generating this diagnostic accordingly. I agree that can be a separate patch since it is orthogonal to this PR.


class KernelFunctor {
public:
void operator()() const {
Expand All @@ -48,6 +51,7 @@ class KernelFunctor {
func5();
func6();
func7();
func8();
}
};

Expand Down