[Clang] Treat ext_vector_type as a regular type attribute

[jhuber6]

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.

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang

Author: Joseph Huber (jhuber6)

Changes

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.

---

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

6 Files Affected:

• (modified) clang/docs/ReleaseNotes.rst (+1)
• (modified) clang/include/clang/Basic/Attr.td (+3-10)
• (modified) clang/include/clang/Basic/AttrDocs.td (+23)
• (modified) clang/lib/Sema/SemaDeclAttr.cpp (+2-1)
• (modified) clang/test/CodeGenCUDA/amdgpu-bf16.cu (+4-8)
• (modified) clang/test/Sema/types.c (+1-1)

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;
 

[AlexVlx]

I'm not super thrilled about NOT having to introduce a type name, I'm not looking forward to the exciting world of void foo(int __attribute__((ext_vector_type(4))) x) replacing void foo(int4 x).

I am extremely not at ease with flipping this to be a C++ attribute. Back in the dawn of time, C++ attributes were meant to be non-semantic / removable with no effect. This has kindof become muddled throughout the years, and we're probably moving away from that, and yes you have namespaced it, but I'd like such things to not be super cute and to remain clearly double prefix scary things.

[jhuber6]

I'm not super thrilled about NOT having to introduce a type name, I'm not looking forward to the exciting world of void foo(int __attribute__((ext_vector_type(4))) x) replacing void foo(int4 x).

This makes introducing a type name easier because you can use using with templates. But I really don't think that it makes sense to restrict this to only typedef, it's even listed as a FIXME in the original patch.

I am extremely not at ease with flipping this to be a C++ attribute. Back in the dawn of time, C++ attributes were meant to be non-semantic / removable with no effect. This has kindof become muddled throughout the years, and we're probably moving away from that, and yes you have namespaced it, but I'd like such things to not be super cute and to remain clearly double prefix scary things.

I think 301eb6b tried to solve some of those issues. A lot of type / expression level attributes are expected to be droppable, but we have tons of existing ones that definitely aren't. E.g. address spaces and the fifty different vector types for each target. I could force it to be double prefixed, but this really still should be a type level attribute, since it maps 1-to-1 with the LLVM vector type. I'd greatly prefer not to arbitrarily restrict it though, since it would be a major outlier.

[jhuber6]

I suppose I could try to introduce the RegularKeyword usage that the other patch used? That would make it like __ext_vector_type I think. (And honestly, while I'm at it, do you think we could drop the ext?)

[AlexVlx]

I'm not super thrilled about NOT having to introduce a type name, I'm not looking forward to the exciting world of void foo(int __attribute__((ext_vector_type(4))) x) replacing void foo(int4 x).

This makes introducing a type name easier because you can use using with templates. But I really don't think that it makes sense to restrict this to only typedef, it's even listed as a FIXME in the original patch.

It is not restricted to typedef today, you can use it with using: https://gcc.godbolt.org/z/j8xaqcMTo. Not sure what you have in mind when you mention templates, there were indeed some oddities around them in dependent contexts, but this patch wouldn't address those I don't think so it's possible they got sorted out. Could you please clarify?

I could force it to be double prefixed, but this really still should be a type level attribute, since it maps 1-to-1 with the LLVM vector type. I'd greatly prefer not to arbitrarily restrict it though, since it would be a major outlier.

I think you misunderstand, possibly because of my wording. I am not asking you to double prefix, I am saying you should not make this a C++ attribute, and leave it exclusively as is is today i.e. do not change it's spelling from GNU to Clang, which was called out in the initial design (confusingly tied to OCL, but different kettle of fish). It is perfectly fine and desirable for this to only be available as __attribute__((ext_vector_type))

[arsenm]

It completely changes the type, it's much more aggressive than a type attribute? sizeof is no longer the same

[jhuber6]

It completely changes the type, it's much more aggressive than a type attribute? sizeof is no longer the same

We already have type-level vector attributes, like the vector size or NEON attributes. Some of them also have clang:: spellings.

[erichkeane]

I don't like the keyword here, particuarly when it doesn't match the attribute. For better or worse, these types are a wart from the GNU variants of these, and we should stick it that way. As a keyword it isn't clear whether this should be treated as a qualifier or specifier, and this patch doesn't answer that at all.

Additionally, the Clang spelling was the right one, this is a clang attribute, we should spell it like 'clang.

[jhuber6]

That was my preference as well. I went ahead and kept the keyword with the ext_vector_type spelling since that should resolve cases where we're worried about double parsing. Should I just keep it as just clang?

[jhuber6]

I don't know why the documentation build keeps failing, just says it's missing a header, but this one looks like all the other ones.

[AaronBallman]

I am extremely not at ease with flipping this to be a C++ attribute. Back in the dawn of time, C++ attributes were meant to be non-semantic / removable with no effect.

This only applies to standard attributes and has never been intended to apply to vendor attributes. Consider other type attributes like calling conventions on functions.

[AaronBallman]

It completely changes the type, it's much more aggressive than a type attribute? sizeof is no longer the same

address_space could be argued to be similarly aggressive, _Atomic as a qualifier specifier also has the same property of not requiring the size and alignment to be the same.

(Note, I'm not saying to do things one way or the other, just pointing out that the type system is filled with oddities already.)

I think vector types as a type attribute are reasonable, at least at first blush. However, if we're talking about redesigning the feature to expose it in a more natural way, we should consider what proposals are in front of the C and C++ standards committees. For example, this is an in-progress proposal that's expected to come in front of WG14 at some point (so the committee has not seen it or weighed in on it): https://www.aerotri.es/C/Array_notation_for_vectorization.pdf

[jhuber6]

I think vector types as a type attribute are reasonable, at least at first blush. However, if we're talking about redesigning the feature to expose it in a more natural way, we should consider what proposals are in front of the C and C++ standards committees.

This mostly just makes these easier to use, they come from https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_C.html#built-in-vector-data-types originally but they map pretty well to the LLVM vectors. I agree it would be nice to have this type of stuff in the standard someday.

[jhuber6]

@erichkeane Is this good to go now? CI's green and I reverted the attribute back to your initial approval.

[erichkeane]

2 test change requests, else LGTM.