[Clang] Treat ext_vector_type as a regular type attribute

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.

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 {

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 = [{

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

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;
 }

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;