[clang] [Clang] Treat `ext_vector_type` as a regular type attribute (PR #130177)

[Joseph Huber via cfe-commits]

https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/130177

>From 728e1bd9cccb56a0acaf5abb35fe64cacc5b4ae9 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 6 Mar 2025 15:08:25 -0600
Subject: [PATCH 1/5] [Clang] Treat `ext_vector_type` as a regular type
 attribute

Summary:
This attribute is mostly borrowed from OpenCL, but is useful in general
for accessing the LLVM vector types. Previously the only way to use it
was through typedefs. This patch changes that to allow use as a regular
type attribute, similar to address spaces.
---
 clang/docs/ReleaseNotes.rst           |  1 +
 clang/include/clang/Basic/Attr.td     | 13 +++----------
 clang/include/clang/Basic/AttrDocs.td | 23 +++++++++++++++++++++++
 clang/lib/Sema/SemaDeclAttr.cpp       |  3 ++-
 clang/test/CodeGenCUDA/amdgpu-bf16.cu | 12 ++++--------
 clang/test/Sema/types.c               |  2 +-
 6 files changed, 34 insertions(+), 20 deletions(-)

diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 86bf836b4a999..695c458b36702 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -145,6 +145,7 @@ Adding [[clang::unsafe_buffer_usage]] attribute to a method definition now turns
 related warnings within the method body.
 
 - The ``no_sanitize`` attribute now accepts both ``gnu`` and ``clang`` names.
+- The ``ext_vector_type(n)`` attribute can now be used as a generic type attribute.
 - Clang now diagnoses use of declaration attributes on void parameters. (#GH108819)
 - Clang now allows ``__attribute__((model("small")))`` and
   ``__attribute__((model("large")))`` on non-TLS globals in x86-64 compilations.
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index dc9b462126125..161a4fe8e0f12 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1721,17 +1721,10 @@ def EnableIf : InheritableAttr {
   let Documentation = [EnableIfDocs];
 }
 
-def ExtVectorType : Attr {
-  // This is an OpenCL-related attribute and does not receive a [[]] spelling.
-  let Spellings = [GNU<"ext_vector_type">];
-  // FIXME: This subject list is wrong; this is a type attribute.
-  let Subjects = SubjectList<[TypedefName], ErrorDiag>;
+def ExtVectorType : TypeAttr {
+  let Spellings = [Clang<"ext_vector_type">];
   let Args = [ExprArgument<"NumElements">];
-  let ASTNode = 0;
-  let Documentation = [Undocumented];
-  // This is a type attribute with an incorrect subject list, so should not be
-  // permitted by #pragma clang attribute.
-  let PragmaAttributeSupport = 0;
+  let Documentation = [ExtVectorTypeDocs];
 }
 
 def FallThrough : StmtAttr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index f44fad95423ee..c309b4849b731 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -1113,6 +1113,29 @@ template instantiation, so the value for ``T::number`` is known.
   }];
 }
 
+def ExtVectorTypeDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+
+The ext_vector_type(N) attribute specifies that a type is a vector with N
+elements, directly mapping to an LLVM vector type. Originally from OpenCL, it
+allows element access via [] or x, y, z, w for graphics-style indexing. This
+attribute enables efficient SIMD operations and is usable in general-purpose
+code.
+
+.. code-block:: c++
+
+  template <typename T, uint32_t N>
+  constexpr T simd_reduce(T [[clang::ext_vector_type(N)]] v) {
+    T sum{};
+    for (uint32_t i = 0; i < N; ++i) {
+      sum += v[i];
+    }
+    return sum;
+  }
+  }];
+}
+
 def DiagnoseIfDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 1405ee5341dcf..d32320c581656 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -1191,7 +1191,8 @@ static void handleTestTypestateAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 
 static void handleExtVectorTypeAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   // Remember this typedef decl, we will need it later for diagnostics.
-  S.ExtVectorDecls.push_back(cast<TypedefNameDecl>(D));
+  if (isa<TypedefNameDecl>(D))
+    S.ExtVectorDecls.push_back(cast<TypedefNameDecl>(D));
 }
 
 static void handlePackedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
diff --git a/clang/test/CodeGenCUDA/amdgpu-bf16.cu b/clang/test/CodeGenCUDA/amdgpu-bf16.cu
index 4610b4ae3cbe5..f6533d7faf296 100644
--- a/clang/test/CodeGenCUDA/amdgpu-bf16.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-bf16.cu
@@ -111,19 +111,15 @@ __device__ __bf16 test_call( __bf16 in) {
 // CHECK-NEXT:    ret void
 //
 __device__ void test_vec_assign() {
-  typedef __attribute__((ext_vector_type(2))) __bf16 bf16_x2;
-  bf16_x2 vec2_a, vec2_b;
+  __bf16 [[clang::ext_vector_type(2)]] vec2_a, vec2_b;
   vec2_a = vec2_b;
 
-  typedef __attribute__((ext_vector_type(4))) __bf16 bf16_x4;
-  bf16_x4 vec4_a, vec4_b;
+  __bf16 [[clang::ext_vector_type(4)]] vec4_a, vec4_b;
   vec4_a = vec4_b;
 
-  typedef __attribute__((ext_vector_type(8))) __bf16 bf16_x8;
-  bf16_x8 vec8_a, vec8_b;
+  __bf16 [[clang::ext_vector_type(8)]] vec8_a, vec8_b;
   vec8_a = vec8_b;
 
-  typedef __attribute__((ext_vector_type(16))) __bf16 bf16_x16;
-  bf16_x16 vec16_a, vec16_b;
+  __bf16 [[clang::ext_vector_type(16)]] vec16_a, vec16_b;
   vec16_a = vec16_b;
 }
diff --git a/clang/test/Sema/types.c b/clang/test/Sema/types.c
index 2a5f530740e9a..2be0e6544f3d7 100644
--- a/clang/test/Sema/types.c
+++ b/clang/test/Sema/types.c
@@ -78,7 +78,7 @@ typedef int __attribute__((ext_vector_type(0))) e4;                // expected-e
 // no support for vector enum type
 enum { e_2 } x3 __attribute__((vector_size(64))); // expected-error {{invalid vector element type}}
 
-int x4 __attribute__((ext_vector_type(64)));  // expected-error {{'ext_vector_type' attribute only applies to typedefs}}
+int x4 __attribute__((ext_vector_type(64)));
 
 typedef __attribute__ ((ext_vector_type(32),__aligned__(32))) unsigned char uchar32;
 

>From 74ee49f1991aca82746776ed515daa3a19b3efb2 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 6 Mar 2025 16:37:52 -0600
Subject: [PATCH 2/5] more tests

---
 clang/test/CodeGen/vector.c  |  7 ++++---
 clang/test/Sema/vector-ast.c | 13 +++++++++++++
 2 files changed, 17 insertions(+), 3 deletions(-)
 create mode 100644 clang/test/Sema/vector-ast.c

diff --git a/clang/test/CodeGen/vector.c b/clang/test/CodeGen/vector.c
index c20e320463056..5d677aaf16948 100644
--- a/clang/test/CodeGen/vector.c
+++ b/clang/test/CodeGen/vector.c
@@ -8,6 +8,10 @@ void test1(void) {
 __v4hi x = {1,2,3};
 __v4hi y = {1,2,3,4};
 
+
+// CHECK: @z = local_unnamed_addr global <8 x float> zeroinitializer
+float z __attribute__((ext_vector_type(8)));
+
 typedef int vty __attribute((vector_size(16)));
 int test2(void) { vty b; return b[2LL]; }
 
@@ -18,9 +22,6 @@ void test3 ( vec4* a, char b, float c ) {
   (*a)[b] = c;
 }
 
-
-
-
 #include <mmintrin.h>
 
 int test4(int argc, char *argv[]) {
diff --git a/clang/test/Sema/vector-ast.c b/clang/test/Sema/vector-ast.c
new file mode 100644
index 0000000000000..4293055c7c97a
--- /dev/null
+++ b/clang/test/Sema/vector-ast.c
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 %s -verify -ast-dump | FileCheck %s
+
+// expected-no-diagnostics
+
+// CHECK: VarDecl {{.*}} x 'int __attribute__((ext_vector_type(4)))'
+int x __attribute__((ext_vector_type(4)));
+
+// CHECK: FunctionDecl {{.*}} 'int () __attribute__((ext_vector_type(4)))'
+int __attribute__((ext_vector_type(4))) foo() { return x; }
+// CHECK:  CompoundStmt
+// CHECK-NEXT: ReturnStmt
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int __attribute__((ext_vector_type(4)))' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int __attribute__((ext_vector_type(4)))' lvalue Var {{.*}} 'x' 'int __attribute__((ext_vector_type(4)))'

>From abe0608c409f8dc7eb58df1c2c8f0b91174c0287 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 6 Mar 2025 16:56:12 -0600
Subject: [PATCH 3/5] better docs

---
 clang/include/clang/Basic/AttrDocs.td | 15 ++++++++++++---
 1 file changed, 12 insertions(+), 3 deletions(-)

diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index c309b4849b731..a245614e7da43 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -1119,9 +1119,10 @@ def ExtVectorTypeDocs : Documentation {
 
 The ext_vector_type(N) attribute specifies that a type is a vector with N
 elements, directly mapping to an LLVM vector type. Originally from OpenCL, it
-allows element access via [] or x, y, z, w for graphics-style indexing. This
-attribute enables efficient SIMD operations and is usable in general-purpose
-code.
+allows element access the array subscript operator ``[]``, ``sN`` where ``N`` is
+a hexadecimal value, or ``x``, ``y``, ``z``, ``w`` for graphics-style indexing.
+This attribute enables efficient SIMD operations and is usable in
+general-purpose code.
 
 .. code-block:: c++
 
@@ -1133,6 +1134,14 @@ code.
     }
     return sum;
   }
+
+The vector type also supports swizzling up to sixteen elements. This can be done
+using the object accessors.
+.. code-block:: c++
+
+  using f16_x16 = _Float16 __attribute__((ext_vector_type(16)));
+
+  f16_x16 reverse(f16_x16 v) { return v.sfedcba9876543210; }
   }];
 }
 

>From 7a9233fa3dca3672f9a76d0f4b9bbe94caf63697 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 6 Mar 2025 17:14:41 -0600
Subject: [PATCH 4/5] reduction thing

---
 clang/include/clang/Basic/AttrDocs.td | 13 ++++++++-----
 1 file changed, 8 insertions(+), 5 deletions(-)

diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index a245614e7da43..062fcef475a30 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -1128,15 +1128,18 @@ general-purpose code.
 
   template <typename T, uint32_t N>
   constexpr T simd_reduce(T [[clang::ext_vector_type(N)]] v) {
-    T sum{};
-    for (uint32_t i = 0; i < N; ++i) {
-      sum += v[i];
+    static_assert((N & (N - 1)) == 0, "N must be a power of two");
+    if constexpr (N == 1) {
+      return v[0];
+    } else {
+      T [[clang::ext_vector_type(N / 2)]] reduced = v.hi + v.lo;
+      return simd_reduce<T, N / 2>(reduced);
     }
-    return sum;
   }
 
 The vector type also supports swizzling up to sixteen elements. This can be done
-using the object accessors.
+using the object accessors. The OpenCL documentation lists the full list of
+accepted values.
 .. code-block:: c++
 
   using f16_x16 = _Float16 __attribute__((ext_vector_type(16)));

>From 1a9bf54fddb1de2797de405182fcc228f904debf Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 6 Mar 2025 19:28:26 -0600
Subject: [PATCH 5/5] Remove CXX attr

---
 clang/include/clang/Basic/Attr.td     | 2 +-
 clang/include/clang/Basic/AttrDocs.td | 6 +++---
 clang/test/CodeGenCUDA/amdgpu-bf16.cu | 8 ++++----
 3 files changed, 8 insertions(+), 8 deletions(-)

diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 161a4fe8e0f12..36fe2eaf6fb31 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1722,7 +1722,7 @@ def EnableIf : InheritableAttr {
 }
 
 def ExtVectorType : TypeAttr {
-  let Spellings = [Clang<"ext_vector_type">];
+  let Spellings = [GNU<"ext_vector_type">];
   let Args = [ExprArgument<"NumElements">];
   let Documentation = [ExtVectorTypeDocs];
 }
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 062fcef475a30..43c5dc0351256 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -1127,13 +1127,13 @@ general-purpose code.
 .. code-block:: c++
 
   template <typename T, uint32_t N>
-  constexpr T simd_reduce(T [[clang::ext_vector_type(N)]] v) {
+  constexpr T simd_reduce(T __attribute__((ext_vector_type(N))) v) {
     static_assert((N & (N - 1)) == 0, "N must be a power of two");
     if constexpr (N == 1) {
       return v[0];
     } else {
-      T [[clang::ext_vector_type(N / 2)]] reduced = v.hi + v.lo;
-      return simd_reduce<T, N / 2>(reduced);
+      T __attribute__((ext_vector_type(N / 2))) reduced = v.hi + v.lo;
+      return simd_reduce(reduced);
     }
   }
 
diff --git a/clang/test/CodeGenCUDA/amdgpu-bf16.cu b/clang/test/CodeGenCUDA/amdgpu-bf16.cu
index f6533d7faf296..027ed4a006014 100644
--- a/clang/test/CodeGenCUDA/amdgpu-bf16.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-bf16.cu
@@ -111,15 +111,15 @@ __device__ __bf16 test_call( __bf16 in) {
 // CHECK-NEXT:    ret void
 //
 __device__ void test_vec_assign() {
-  __bf16 [[clang::ext_vector_type(2)]] vec2_a, vec2_b;
+  __bf16 __attribute__((ext_vector_type(2))) vec2_a, vec2_b;
   vec2_a = vec2_b;
 
-  __bf16 [[clang::ext_vector_type(4)]] vec4_a, vec4_b;
+  __bf16 __attribute__((ext_vector_type(4))) vec4_a, vec4_b;
   vec4_a = vec4_b;
 
-  __bf16 [[clang::ext_vector_type(8)]] vec8_a, vec8_b;
+  __bf16 __attribute__((ext_vector_type(8))) vec8_a, vec8_b;
   vec8_a = vec8_b;
 
-  __bf16 [[clang::ext_vector_type(16)]] vec16_a, vec16_b;
+  __bf16 __attribute__((ext_vector_type(16))) vec16_a, vec16_b;
   vec16_a = vec16_b;
 }