[clang] [CLANG][CUDA] Enable alignas after GNU attributes (PR #133105)

[via cfe-commits]

https://github.com/DenisGZM created https://github.com/llvm/llvm-project/pull/133105

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.

>From 54491421730f73cd25d36cf2aaadef3fddd2220e Mon Sep 17 00:00:00 2001
From: Denis Gerasimov <Denis.Gerasimov at baikalelectronics.ru>
Date: Wed, 26 Mar 2025 18:08:26 +0300
Subject: [PATCH] [CLANG] Enable alignas after GNU attributes

---
 clang/lib/Parse/ParseStmt.cpp          |  5 +++++
 clang/test/SemaCUDA/cuda-attr-order.cu | 15 +++++++++++++++
 2 files changed, 20 insertions(+)
 create mode 100644 clang/test/SemaCUDA/cuda-attr-order.cu

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..e9b79fbfdba9a
--- /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];
+}
\ No newline at end of file