Skip to content

[SYCL] Update kernel name diagnostics logic #3069

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

Closed
srividya-sundaram wants to merge 4 commits into from
Closed

[SYCL] Update kernel name diagnostics logic #3069

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b96891f
Allow nested structs as kernel names
srividya-sundaram Jan 16, 2021
55efe70
[SYCL] Update kernel name diagnostics logic
srividya-sundaram Jan 20, 2021
f7fe8cc
Update test to remove anonymous namespace functor
srividya-sundaram Jan 21, 2021
a71a016
clang-format fix
srividya-sundaram Jan 21, 2021
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
50 changes: 35 additions & 15 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3025,19 +3025,34 @@ class SYCLKernelNameTypeVisitor
S.getASTContext().getLangOpts().SYCLUnnamedLambda;
const DeclContext *DeclCtx = Tag->getDeclContext();
if (DeclCtx && !UnnamedLambdaEnabled) {
auto *NameSpace = dyn_cast_or_null<NamespaceDecl>(DeclCtx);
if (NameSpace && NameSpace->isStdNamespace()) {
S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
<< KernelNameType;
S.Diag(KernelInvocationFuncLoc, diag::note_invalid_type_in_sycl_kernel)
<< /* kernel name cannot be a type in the std namespace */ 2
<< QualType(Tag->getTypeForDecl(), 0);
IsInvalid = true;
return;
}
if (!DeclCtx->isTranslationUnit() && !isa<NamespaceDecl>(DeclCtx)) {
const bool KernelNameIsMissing = Tag->getName().empty();
if (KernelNameIsMissing) {

while (!DeclCtx->isTranslationUnit()) {
auto *NSDecl = dyn_cast_or_null<NamespaceDecl>(DeclCtx);
Copy link
Contributor

@erichkeane erichkeane Jan 21, 2021

Choose a reason for hiding this comment

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

This can just be a dyn_cast, otherwise the condition in the 'while' is UB :)

if (NSDecl && NSDecl->isStdNamespace()) {
Copy link
Contributor

@elizabethandrews elizabethandrews Jan 21, 2021

Choose a reason for hiding this comment

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

Recursive check will now throw an error for std::test::U, which previously passed. I think the error is right behavior though.

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021

Choose a reason for hiding this comment

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

You're right that std::test::U should fail, however the loop makes me concerned that the visitor isn't doing what it is supposed to. Why is it not doing a 'visit' of the namespace and catching it that way? I would expect each of the decl-contexts to be visited, not need to be looped here.

Copy link
Contributor

@Fznamznon Fznamznon Jan 22, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

This is a TypeVisitor and ConstTemplateArgumentVisitor, it can visit only types and template arguments, we have to inherit it from some "declaration context visitor" to do so, if this thing exists at all.

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021

Choose a reason for hiding this comment

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

Ah, huh... I figured we would use something like the text-dumpers use, which visit each node along the way individually. They seem to better fit the integration-header/kernel name checking.

Copy link
Contributor

@Fznamznon Fznamznon Jan 22, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

Do you mean things like TextNodeDumper? i.e. this one https://clang.llvm.org/doxygen/classclang_1_1TextNodeDumper.html .
The original idea was to make our visitor similar to this one. That is why we used TypeVisitor and ConstTemplateArgumentVisitor as bases, as well as TextNodeDumper does. We could take a look on how it handles namespaces and other declaration contexts, though.

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021

Choose a reason for hiding this comment

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

Thanks for the context.

We definitely need to prevent std::something::SomethingElse, the fact that we aren't is likely a regression (as I seem to remember it worked before the refactor? Unfortunate there were no tests).

This patch does NOT seem to do anything with the integration header though, which would need to be taught how to forward-declare these globally visible names. That said, I don't think this necessary part of this is actually possible with integration-headers (see my other comment).

At the moment, I'd suggest a separate patch to fix the std::something::something_else case (as well as the something::<anon_ns>::another::type case), and we wait for the nested-struct version until we can get @mkinsner and the SYCL spec clarifications of what we have to do.

if we have to support this nested struct case, I think we have to figure out an alternative to integration headers for runtime support.

Copy link
Contributor Author

@srividya-sundaram srividya-sundaram Jan 22, 2021

Choose a reason for hiding this comment

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

as I seem to remember it worked before the refactor? Unfortunate there were no tests)

Refactor did not break anything. std::something::SomethingElse was not diagnosed before.
Same with anonymous NS.

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

Interesting. That diffs from my memory(though it is historically not particularly accurate, so I can buy I mis-remembered), but still needs fixing. The rule to disallow std::string has the same motivation/conclusion as disallowing std::ranges::views::view_base

Copy link
Contributor Author

@srividya-sundaram srividya-sundaram Jan 22, 2021

Choose a reason for hiding this comment

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

You can look at the history. I can fix it, just clarifying that it was something not previously supported and something I encountered while working on the original issue.

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021

Choose a reason for hiding this comment

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

I believe you :) I could swear it was caught at SOME point (as I was sure I encountered it at one point...), but it doesn't mean it survived until the refactor in question I guess.

S.Diag(KernelInvocationFuncLoc,
diag::err_sycl_kernel_incorrectly_named)
<< KernelNameType;
S.Diag(KernelInvocationFuncLoc,
diag::note_invalid_type_in_sycl_kernel)
<< /* kernel name cannot be a type in the std namespace */ 2
<< QualType(Tag->getTypeForDecl(), 0);
IsInvalid = true;
return;
}
if (NSDecl && NSDecl->isAnonymousNamespace()) {
S.Diag(KernelInvocationFuncLoc,
diag::err_sycl_kernel_incorrectly_named)
<< KernelNameType;
S.Diag(KernelInvocationFuncLoc,
diag::note_invalid_type_in_sycl_kernel)
<< /* kernel name is not globally-visible */ 0
<< QualType(Tag->getTypeForDecl(), 0);
IsInvalid = true;
return;
}

const bool UnnamedTypeUsed = Tag->getName().empty();
Copy link
Contributor

@elizabethandrews elizabethandrews Jan 21, 2021

Choose a reason for hiding this comment

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

This can be outside the while loop right?

if (UnnamedTypeUsed) {
S.Diag(KernelInvocationFuncLoc,
diag::err_sycl_kernel_incorrectly_named)
<< KernelNameType;
Expand All @@ -3047,7 +3062,7 @@ class SYCLKernelNameTypeVisitor
IsInvalid = true;
return;
}
if (Tag->isCompleteDefinition()) {
if (isa<FunctionDecl>(DeclCtx) && Tag->isCompleteDefinition()) {
Copy link
Contributor Author

@srividya-sundaram srividya-sundaram Jan 21, 2021

Choose a reason for hiding this comment

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

The if at L3065 is to handle cases like the following:

       void test() {
        class InvalidKernelName1 {};
          q.submit([&](handler &h) {
            h.single_task<InvalidKernelName1>([]() {});
          });

The order of decl context for InvalidKernelName1 is
clang::Decl::Function
clang::Decl::TranslationUnit

And class InvalidKernelName1 {}; is completely defined.

Copy link
Contributor

@erichkeane erichkeane Jan 21, 2021

Choose a reason for hiding this comment

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

I guess I'm more asking for a situation where that isa is required. It seems that both your examples are positive things here, but what cases are you trying to exclude from these 'if's?

Copy link
Contributor Author

@srividya-sundaram srividya-sundaram Jan 21, 2021

Choose a reason for hiding this comment

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

Trying to exclude cases like nested named structs and classes inside structs.

Copy link
Contributor

@elizabethandrews elizabethandrews Jan 21, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

Hmm I see... You're limiting the check to only function scope. I think it might be better to instead check if the scope is file context or something, and throw the error if it isn't. @erichkeane thoughts?

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021

Choose a reason for hiding this comment

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

I unfortunately don't know the visitor here well enough to know exactly what is going on here, but this part of the change "feels" wrong, but I obviously cannot specify what makes me think that. It just generally seems like the visitor isn't doing what it is supposed to if all this is necessary, and these two 'if' statements are particularly constrained for situations that are not sufficiently thought out.

I'm hoping that @Fznamznon can analyze the patch/problem and make sure this is the right approach.

Copy link
Contributor

@Fznamznon Fznamznon Jan 22, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

So, the original purpose of this PR was to fix cases when error was reported for nested name specifiers, i.e.

  struct A {
    struct B {
    };
  };
  q.submit([&](handler &h) { h.single_task<A::B>([]() {});

Turns out it is a valid case and we diagnosed that.
The spec rule which we are trying to satisfy here is "kernel name should be globally-visible".
I believe I don't know exactly what it means, but I asked Sri to clarify it before preparing a patch. My gut feeling tells that class name defined inside function scope is not globally-visible, so when we discussed this patch we decided to diagnose if kernel name was defined in function/method scope. Seems to me if my gut tells us right thing, we need to throw "kernel name is not globally-visible" error if we encounter a function/method scope during our decl context loop. I'm not really sure, why we need to check on complete definition here at all. And why we are checking context on line 3076

Copy link
Contributor

@erichkeane erichkeane Jan 22, 2021

Choose a reason for hiding this comment

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

You're correct in a class inside a function scope is not globally visible. So when the 'visitor' finds a function-type/decl, I would assume it would reject it.

Globally visible names appear only in namespaces, record types, or enums I believe (though I don't believe you can create a record-type inside an enum in any way, unless there is an incredible trick I'm missing).

Copy link
Contributor Author

@srividya-sundaram srividya-sundaram Jan 22, 2021 •
edited
Loading

Choose a reason for hiding this comment

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

I believe I don't know exactly what it means, but I asked Sri to clarify it before preparing a patch.

I looked at "5.2 Naming of kernels" in the SYCL2020-provisional-spec and based on my understanding of rule#2,3,4 made this change.

S.Diag(KernelInvocationFuncLoc,
diag::err_sycl_kernel_incorrectly_named)
<< KernelNameType;
Expand All @@ -3056,11 +3071,16 @@ class SYCLKernelNameTypeVisitor
<< /* kernel name is not globally-visible */ 0
<< QualType(Tag->getTypeForDecl(), 0);
IsInvalid = true;
} else {
return;
}
if (isa<CXXMethodDecl>(DeclCtx) && !Tag->isCompleteDefinition()) {
Copy link
Contributor

@erichkeane erichkeane Jan 21, 2021

Choose a reason for hiding this comment

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

This and line 3065 seems to change the logic quite a bit here. Can you explain what is happening here?

Copy link
Contributor Author

@srividya-sundaram srividya-sundaram Jan 21, 2021

Choose a reason for hiding this comment

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

The if block at L3076 is to handle the cases like:

        int main()
        h.single_task<class fake_kernel2>([]() {
              auto l = [](auto f) { f(); };
        });

The order of DeclContexts for fake_kernel2 is
clang::Decl::CXXMethod
clang::Decl::CXXRecord
clang::Decl::Function
clang::Decl::TranslationUnit

So I just check if the DeclCtx is clang::Decl::CXXMethod and "class fake_kernel2" does not have complete definition.

Copy link
Contributor

@elizabethandrews elizabethandrews Jan 21, 2021

Choose a reason for hiding this comment

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

Do you need the context check here? Won't the warning for implicit decl be generated even without this check?

S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl);
S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl)
<< Tag->getName();
}

// Repeat the above checks for DeclCtx's in the parent-chain
DeclCtx = DeclCtx->getParent();
}
}
}
Expand Down
12 changes: 2 additions & 10 deletions clang/test/CodeGenSYCL/int_header1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s -o %t.out
// RUN: %clang_cc1 -fsycl -fsycl-is-device -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out
// RUN: FileCheck -input-file=%t.h %s

// CHECK:template <> struct KernelInfo<class KernelName> {
Expand All @@ -11,7 +11,6 @@
// CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName8<::nm1::nm2::C>> {
// CHECK:template <> struct KernelInfo<::TmplClassInAnonNS<ClassInAnonNS>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName9<char>> {
// CHECK:template <> struct KernelInfo<::nm1::KernelName3<const volatile ::nm1::KernelName3<const volatile char>>> {

Expand Down Expand Up @@ -49,11 +48,6 @@ namespace nm1 {

} // namespace nm1

namespace {
class ClassInAnonNS;
template <typename T> class TmplClassInAnonNS;
}

struct MyWrapper {
class KN101 {};

Expand Down Expand Up @@ -130,8 +124,6 @@ struct MyWrapper {

// kernel name type is a templated class, both the top-level class and the
// template argument are declared in the anonymous namespace
kernel_single_task<TmplClassInAnonNS<class ClassInAnonNS>>(
[=]() { acc.use(); });

// Kernel name type is a templated specialization class with empty template pack argument
kernel_single_task<nm1::KernelName9<char>>(
Expand Down Expand Up @@ -165,7 +157,7 @@ int main() {
KernelInfo<class nm1::KernelName3<class KernelName5>>::getName();
KernelInfo<class nm1::KernelName4<class KernelName7>>::getName();
KernelInfo<class nm1::KernelName8<nm1::nm2::C>>::getName();
KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>>::getName();

KernelInfo<class nm1::KernelName9<char>>::getName();
#endif //__SYCL_DEVICE_ONLY__
}
69 changes: 69 additions & 0 deletions clang/test/SemaSYCL/allow-nested-structs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,69 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s

#include "sycl.hpp"

struct NestedStruct1 {
struct NestedStruct2 {
struct NestedStruct3 {};
};
};

namespace {
struct StructInAnonymousNS {};
} // namespace

namespace ValidNS {
struct StructinValidNS {};
} // namespace ValidNS

struct Parent {
using A = struct {
struct Child1 {
struct Child2 {};
};
};
};

struct MyWrapper {

public:
void test() {
cl::sycl::queue q;
struct StructInsideFunc {};

// no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<NestedStruct1::NestedStruct2::NestedStruct3>([] {});
});

// no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<ValidNS::StructinValidNS>([] {});
});

// no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<Parent::A::Child1::Child2>([] {});
});

// expected-error@Inputs/sycl.hpp:220 {{'(anonymous namespace)::StructInAnonymousNS' is an invalid kernel name type}}
// expected-note@Inputs/sycl.hpp:220 {{'(anonymous namespace)::StructInAnonymousNS' should be globally-visible}}
// expected-note@+2{{in instantiation of function template specialization}}
q.submit([&](cl::sycl::handler &h) {
h.single_task<StructInAnonymousNS>([] {});
});

// expected-error@Inputs/sycl.hpp:220 {{'StructInsideFunc' is an invalid kernel name type}}
// expected-note@Inputs/sycl.hpp:220 {{'StructInsideFunc' should be globally-visible}}
// expected-note@+2{{in instantiation of function template specialization}}
q.submit([&](cl::sycl::handler &h) {
h.single_task<StructInsideFunc>([] {});
});
}
};

int main() {
cl::sycl::queue q;

return 0;
}
26 changes: 5 additions & 21 deletions clang/test/SemaSYCL/unnamed-kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,20 +52,12 @@ struct MyWrapper {
h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}}
// expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' should be globally-visible}}
// expected-note@+3{{in instantiation of function template specialization}}
#endif
// no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidKernelName0>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' is an invalid kernel name type}}
// expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName3' should be globally-visible}}
// expected-note@+3{{in instantiation of function template specialization}}
#endif
//no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
});
Expand All @@ -85,21 +77,13 @@ struct MyWrapper {
});

using InvalidAlias = InvalidKernelName4;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}}
// expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' should be globally-visible}}
// expected-note@+3{{in instantiation of function template specialization}}
#endif
// no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidAlias>([] {});
});

using InvalidAlias1 = InvalidKernelName5;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' is an invalid kernel name type}}
// expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName5' should be globally-visible}}
// expected-note@+3{{in instantiation of function template specialization}}
#endif
// no error
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
});
Expand All @@ -124,4 +108,4 @@ int main() {
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });

return 0;
}
}
Copy link
Contributor

@keryell keryell Jan 21, 2021

Choose a reason for hiding this comment

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

Restore it.

33 changes: 3 additions & 30 deletions sycl/test/functor/kernel_functor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,27 +18,6 @@ constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;

// Case 1:
// - functor class is defined in an anonymous namespace
// - the '()' operator:
// * does not have parameters (to be used in 'single_task').
// * has the 'const' qualifier
namespace {
class Functor1 {
public:
Functor1(
int X_,
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_)
: X(X_), Acc(Acc_) {}

void operator()() const { Acc[0] += X; }

private:
int X;
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
};
}

// Case 2:
// - functor class is defined in a namespace
// - the '()' operator:
// * does not have parameters (to be used in 'single_task').
Expand All @@ -60,7 +39,7 @@ class Functor2 {
};
}

// Case 3:
// Case 2:
// - functor class is templated and defined in the translation unit scope
// - the '()' operator:
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
Expand All @@ -78,7 +57,7 @@ template <typename T> class TmplFunctor {
cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> Acc;
};

// Case 4:
// Case 3:
// - functor class is templated and defined in the translation unit scope
// - the '()' operator:
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
Expand All @@ -103,12 +82,6 @@ int foo(int X) {
cl::sycl::queue Q;
cl::sycl::buffer<int, 1> Buf(A, 1);

Q.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
Functor1 F(X, Acc);

cgh.single_task(F);
});
Q.submit([&](cl::sycl::handler &cgh) {
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
ns::Functor2 F(X, Acc);
Expand Down Expand Up @@ -167,7 +140,7 @@ template <typename T> T bar(T X) {
int main() {
const int Res1 = foo(10);
const int Res2 = bar(10);
const int Gold1 = 40;
const int Gold1 = 30;
const int Gold2 = 80;

assert(Res1 == Gold1);
Expand Down