[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,42 @@ 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 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++
+
+  template <typename T, uint32_t N>
+  constexpr T simd_reduce(T [[clang::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(reduced);
+    }
+  }
+
+The vector type also supports swizzling up to sixteen elements. This can be done
+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)));
+
+  f16_x16 reverse(f16_x16 v) { return v.sfedcba9876543210; }
+
+See the OpenCL documentation for some more complete examples.
+  }];
+}
+
 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/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[]) {

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;
 

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