[SYCL] Update kernel name diagnostics logic #3069

srividya-sundaram · 2021-01-21T16:10:14Z

This patch makes the following changes:

Allows nested named types to be used as valid kernel names.
Allows kernel names inside named namespaces.
Disallows kernel names inside anonymous namespaces.

It does the above by recursively checking if the declaration contexts in the parent-chain of the kernel name are valid or not.

PS: The standards quote (in C++20 as [class.access]p5), clarifies that visibility isn't controlled by access specifiers (public/private/etc).
Before this patch, the compiler would throw an error if any of the InvalidKernelName* was used as kernel name :

struct MyWrapper {
private:
  class InvalidKernelName0 {};
  class InvalidKernelName3 {};
  class InvalidKernelName4 {};
  class InvalidKernelName5 {};
}

But based on standards quote, access specifiers do not matter and we should allow nested named types.

vladimirlaz

LGTM for test changes
But you should also fix llvm-test-suite test (SYCL :: Functor/kernel_functor.cpp):
http://icl-jenkins.sc.intel.com:8080/job/SYCL_CI/job/intel/job/Win/job/Test_Suite/630/console

erichkeane · 2021-01-21T19:41:26Z

clang/lib/Sema/SemaSYCL.cpp

-        if (KernelNameIsMissing) {
+
+      while (!DeclCtx->isTranslationUnit()) {
+        auto *NSDecl = dyn_cast_or_null<NamespaceDecl>(DeclCtx);


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

srividya-sundaram · 2021-01-21T19:41:55Z

LGTM for test changes
But you should also fix llvm-test-suite test (SYCL :: Functor/kernel_functor.cpp):
http://icl-jenkins.sc.intel.com:8080/job/SYCL_CI/job/intel/job/Win/job/Test_Suite/630/console

I addressed SYCL :: Functor/kernel_functor.cpp in this commit:

f7fe8cc

@vladimirlaz I don't see any errors shown in the log for the 2 CUDA failing tests. I just see the RUN command and "Skipping test" in the log. Do you have any idea what's going on?

erichkeane · 2021-01-21T19:42:38Z

clang/lib/Sema/SemaSYCL.cpp

-        } else {
+          return;
+        }
+        if (isa<CXXMethodDecl>(DeclCtx) && !Tag->isCompleteDefinition()) {


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

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.

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

srividya-sundaram · 2021-01-21T19:53:28Z

clang/lib/Sema/SemaSYCL.cpp

@@ -3047,7 +3062,7 @@ class SYCLKernelNameTypeVisitor
          IsInvalid = true;
          return;
        }
-        if (Tag->isCompleteDefinition()) {
+        if (isa<FunctionDecl>(DeclCtx) && Tag->isCompleteDefinition()) {


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.

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?

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

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?

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.

keryell

Nice to have more tests and diagnostics.

keryell · 2021-01-21T22:40:08Z

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -124,4 +108,4 @@ int main() {
  q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });

  return 0;
-}
+}


Restore it.

elizabethandrews · 2021-01-21T23:32:14Z

clang/lib/Sema/SemaSYCL.cpp

+          return;
+        }
+
+        const bool UnnamedTypeUsed = Tag->getName().empty();


This can be outside the while loop right?

elizabethandrews · 2021-01-21T23:51:48Z

clang/lib/Sema/SemaSYCL.cpp

-        } else {
+          return;
+        }
+        if (isa<CXXMethodDecl>(DeclCtx) && !Tag->isCompleteDefinition()) {


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

elizabethandrews · 2021-01-21T23:57:34Z

clang/lib/Sema/SemaSYCL.cpp

@@ -3047,7 +3062,7 @@ class SYCLKernelNameTypeVisitor
          IsInvalid = true;
          return;
        }
-        if (Tag->isCompleteDefinition()) {
+        if (isa<FunctionDecl>(DeclCtx) && Tag->isCompleteDefinition()) {


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?

elizabethandrews · 2021-01-21T23:59:50Z

clang/lib/Sema/SemaSYCL.cpp

+
+      while (!DeclCtx->isTranslationUnit()) {
+        auto *NSDecl = dyn_cast_or_null<NamespaceDecl>(DeclCtx);
+        if (NSDecl && NSDecl->isStdNamespace()) {


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

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.

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.

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.

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.

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.

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.

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

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.

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.

vladimirlaz · 2021-01-22T05:39:05Z

LGTM for test changes
But you should also fix llvm-test-suite test (SYCL :: Functor/kernel_functor.cpp):
http://icl-jenkins.sc.intel.com:8080/job/SYCL_CI/job/intel/job/Win/job/Test_Suite/630/console

I addressed SYCL :: Functor/kernel_functor.cpp in this commit:

f7fe8cc
You have fixed in-tree LIT tests which is compile only. But there is E2E test test in intel/llvm-test-suite repo. See Test DPC toolchain

@vladimirlaz I don't see any errors shown in the log for the 2 CUDA failing tests. I just see the RUN command and "Skipping test" in the log. Do you have any idea what's going on?
If you check test history you will see that it is regression introduced by #2765

Fznamznon · 2021-01-22T14:09:16Z

clang/lib/Sema/SemaSYCL.cpp

@@ -3047,7 +3062,7 @@ class SYCLKernelNameTypeVisitor
          IsInvalid = true;
          return;
        }
-        if (Tag->isCompleteDefinition()) {
+        if (isa<FunctionDecl>(DeclCtx) && Tag->isCompleteDefinition()) {


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

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).

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.

erichkeane · 2021-01-22T14:29:00Z

So as this patch went on, I actually question the original motivation here. My current conclusion is that this is actually a SYCL-standard bug (or at least, we have to vastly change how our integration headers work). At the very least, this is an incomplete patch without the integration header work to make the motivating example usable.

The original justification is:

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

Which DOES pass the globally-visible rule. However, the purpose of the globally-visible rule is to allow us to forward declare the types in the integration-header so that we can pass information to the host compilation.

You'll note however, that we actually CANNOT forward declare 'B' here without defining 'A', which we obviously cannot do.

@mkinsner Is this something you can help @srividya-sundaram with? I think we want this as a defect against the SYCL spec, as I don't think we can properly implement this rule.

mkinsner · 2021-01-25T13:42:38Z

@mkinsner Is this something you can help @srividya-sundaram with? I think we want this as a defect against the SYCL spec, as I don't think we can properly implement this rule.

I believe that an unreleased version of the SYCL spec has already fixed this. I can't describe the fix on GitHub until the spec goes public, but offline we should confirm that you consider the fix sufficient to cover this case.

srividya-sundaram · 2021-01-26T18:46:16Z

Based on offline discussion with @mkinsner, we've confirmed that nested name types are no longer an issue in the upcoming SYCL 2020 release. Closing this review.

srividya-sundaram added 2 commits January 15, 2021 17:04

Allow nested structs as kernel names

b96891f

[SYCL] Update kernel name diagnostics logic

55efe70

srividya-sundaram force-pushed the allow-nested-structs branch from 2b69fb0 to 55efe70 Compare January 21, 2021 16:27

srividya-sundaram changed the title ~~Allow nested structs~~ [SYCL] Update kernel name diagnostics logic Jan 21, 2021

srividya-sundaram added 2 commits January 21, 2021 09:45

Update test to remove anonymous namespace functor

f7fe8cc

clang-format fix

a71a016

srividya-sundaram marked this pull request as ready for review January 21, 2021 18:33

srividya-sundaram requested review from elizabethandrews, premanandrao and a team as code owners January 21, 2021 18:33

srividya-sundaram requested review from vladimirlaz and Fznamznon January 21, 2021 18:33

vladimirlaz approved these changes Jan 21, 2021

View reviewed changes

erichkeane reviewed Jan 21, 2021

View reviewed changes

srividya-sundaram commented Jan 21, 2021

View reviewed changes

keryell suggested changes Jan 21, 2021

View reviewed changes

elizabethandrews reviewed Jan 22, 2021

View reviewed changes

Fznamznon reviewed Jan 22, 2021

View reviewed changes

srividya-sundaram linked an issue Jan 26, 2021 that may be closed by this pull request

[SYCL] Diagnose kernel names in nested "std" and "anonymous" namespace #3103

Closed

srividya-sundaram closed this Jan 26, 2021

[SYCL] Update kernel name diagnostics logic #3069

[SYCL] Update kernel name diagnostics logic #3069

Uh oh!

Conversation

srividya-sundaram commented Jan 21, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

vladimirlaz left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

srividya-sundaram commented Jan 21, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

elizabethandrews Jan 21, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

elizabethandrews Jan 21, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Fznamznon Jan 22, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Fznamznon Jan 22, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

erichkeane Jan 22, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

vladimirlaz commented Jan 22, 2021

Uh oh!

Fznamznon Jan 22, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

srividya-sundaram commented Jan 21, 2021 •

edited

Loading

srividya-sundaram commented Jan 21, 2021 •

edited

Loading

elizabethandrews Jan 21, 2021 •

edited

Loading

elizabethandrews Jan 21, 2021 •

edited

Loading

Fznamznon Jan 22, 2021 •

edited

Loading

Fznamznon Jan 22, 2021 •

edited

Loading

erichkeane Jan 22, 2021 •

edited

Loading

Fznamznon Jan 22, 2021 •

edited

Loading

srividya-sundaram Jan 22, 2021 •

edited

Loading