[CLANG] Enable alignas after GNU attributes

[DenisGZM]

Enable parsing alignas attribute after GNU attributes, before ParseDeclaration

This might be useful for cuda code where shared and other specificators may be mixed with align.

I'd be glad to see if there are any better places or other technique to process this attribute without interrupting current flow of parsing.

[llvmbot]

@llvm/pr-subscribers-clang

Author: Denis.G (DenisGZM)

Changes

Enable parsing alignas attribute after GNU attributes, before ParseDeclaration

This might be useful for cuda code where shared and other specificators may be mixed with align.

I'd be glad to see if there are any better places or other technique to process this attribute without interrupting current flow of parsing.

---

Full diff: https://github.com/llvm/llvm-project/pull/133107.diff

2 Files Affected:

• (modified) clang/lib/Parse/ParseStmt.cpp (+5)
• (added) clang/test/SemaCUDA/cuda-attr-order.cu (+15)

diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp
index 150b2879fc94f..33b9f63bcfa08 100644
--- a/clang/lib/Parse/ParseStmt.cpp
+++ b/clang/lib/Parse/ParseStmt.cpp
@@ -296,6 +296,11 @@ StmtResult Parser::ParseStatementOrDeclarationAfterAttributes(
     goto Retry;
   }
 
+  case tok::kw_alignas: {
+    ParseAlignmentSpecifier(CXX11Attrs);
+    goto Retry;
+  }
+
   case tok::kw_template: {
     SourceLocation DeclEnd;
     ParseTemplateDeclarationOrSpecialization(DeclaratorContext::Block, DeclEnd,
diff --git a/clang/test/SemaCUDA/cuda-attr-order.cu b/clang/test/SemaCUDA/cuda-attr-order.cu
new file mode 100644
index 0000000000000..d3bf5b014d1c6
--- /dev/null
+++ b/clang/test/SemaCUDA/cuda-attr-order.cu
@@ -0,0 +1,15 @@
+// Verify that we can parse a simple CUDA file with different attributes order.
+// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda"  -fsyntax-only -verify %s
+// expected-no-diagnostics
+#include "Inputs/cuda.h"
+
+struct alignas(16) float4 {
+    float x, y, z, w;
+};
+
+__attribute__((device)) float func() {
+    __shared__ alignas(alignof(float4)) float As[4][4];  // Both combinations
+    alignas(alignof(float4)) __shared__  float Bs[4][4]; // must be legal
+
+    return As[0][0] + Bs[0][0];
+}

[Sirraide]

At a glance this does seem like the right place to do this, but this is still missing a release note.

It seems like GCC allows e.g. __attribute__(()) alignas(16) int x in any case, so I don’t see why we shouldn’t allow this too. Can you also add some tests that use __attribute__(()) directly and which aren’t CUDA-specific?

Oh, and can you add solmething like this as a test as well:

struct S { __attribute__((deprecated)) alignas(16) int x; };

[Sirraide]

CC @erichkeane in case there’s a specific reason I’m not aware of as to why we currently don’t allow this.

[DenisGZM]

At a glance this does seem like the right place to do this, but this is still missing a release note.

It seems like GCC allows e.g. __attribute__(()) alignas(16) int x in any case, so I don’t see why we shouldn’t allow this too. Can you also add some tests that use __attribute__(()) directly and which aren’t CUDA-specific?

Oh, and can you add solmething like this as a test as well:

struct S { __attribute__((deprecated)) alignas(16) int x; };

Actually this test doesn't work with this patch...

In this case all attributes are processed in ParseDeclarationSpecifiers, which in my first view was the right place to fix, but has way more complicated logic and easy to break diagnostics.

In ParseDeclarationSpecifiers we parse kw__attributes and other CXX11 Attributes and set bool AttrsLastTime = true to check that last parsed piece was attr. Later this block prohibit attributes with AttrsLastTime = false,

ParseDecl.cpp

    DoneWithDeclSpec:
      if (!AttrsLastTime)
        ProhibitAttributes(attrs);

And AttrsLastTime is always false in declarations of the form: <attributes> <type> <identifier> , because last token we parse is type

Another approach i tried is to add processing alignas-cxx11 just like it is done for C: kw__Alignas and kw_alignas (c23).
Well, it do the parsing but later it skips CXX11 attributes when correcting declaration type (assumed that attributes must have been processed before)

[Sirraide]

Hmm, @erichkeane probably knows where this needs to be parsed then; I might take another look at this myself later (because I’m not sure either off the top of my head), but I’m rather busy today unfortunately...

[DenisGZM]

I added parsing all attributes in ParseCXXClassMemberDeclaration before calling ParseDeclarationSpecifiers and it seems to solve problem, but it also changes annotation ranges for struct and class members

[Sirraide]

Seems reasonable to me, but I’d still like @erichkeane to take a look at this as the attributes code owner

[erichkeane]

Parsing of attributes is admittedly the part I'm least comfortable with here. I would love tests for how this interacts with our __declspec spelling attributes though, and to help determine why we wouldn't parse all 3 together here.

As a followup/future direction for some one, there is perhaps value of a MaybeParseAnyAttributes that does all 3 in a loop.

[erichkeane]

I think this looks reasonable? I would like @AaronBallman to stop by though, he might think of some reason why this isn't right per-grammar.

[AaronBallman]

Thanks for the fix! The changes should come with a release note in clang/docs/ReleaseNotes.rst so users know about the fix.

[AaronBallman]

I think this means we now quietly accept this code: https://godbolt.org/z/eb9MhvTWd

[DenisGZM]

No, we don't. For statements alignas is still illegal:

error: 'alignas' attribute cannot be applied to a statement
    8 |   alignas(int) foo();

[AaronBallman]

It's a bit strange that we're parsing into DeclSpecAttrs for __attribute__ and [] but parsing into DeclAttrs for [[]], can you explain that one a bit?

[DenisGZM]

I followed the same rule as in ParseStatementOrDeclaration where CXX11 Attrs are gathered into DeclAttrs and GNU and MS are gathered into DeclSpecAttrs before ParseStatementOrDeclarationAfterAttributes.

It seems that DeclAttrs is used later to adjust type (e.g. set align). For precise details I still need deep code inspection

[AaronBallman]

Replace with MaybeParseAttributes? (Maybe not if the different arg lists are intentional.)

[AaronBallman]

I actually think the MaybeParseCXX11Attributes() on line 3007 may be incorrect (a pre-existing bug), because we seem to quietly accept: https://godbolt.org/z/s9qbd8jqn but the grammar for member-declaration does not allow an attribute list to appear there: https://eel.is/c++draft/class#nt:member-declaration

[AaronBallman]

I'd like to see a test case like which shows the various combinations of orders we can handle (no need to test struct and class separately as those follow the same code paths).

class C {
  __attribute__(()) [[]] alignas(int) int x;
  __attribute__(()) alignas(int) [[]] int y;
  [[]] __attribute__(()) alignas(int) int z;
  alignas(int) [[]] __attribute__(()) int a;
};

We should also be sure to test this in C23 mode to make sure we get the same behavior in C as well as in C++.