Back to clang

Author: jhuber6

clang/include/clang/Basic/Attr.td

@@ -1722,7 +1722,7 @@ def EnableIf : InheritableAttr {
 }
 
 def ExtVectorType : TypeAttr {
-  let Spellings = [RegularKeyword<"__vector_type">, GNU<"ext_vector_type">];
+  let Spellings = [Clang<"ext_vector_type">];
   let Args = [ExprArgument<"NumElements">];
   let Documentation = [ExtVectorTypeDocs];
 }

clang/include/clang/Basic/AttrDocs.td

@@ -1126,12 +1126,12 @@ general-purpose code.
 .. code-block:: c++
 
   template <typename T, uint32_t N>
-  constexpr T simd_reduce(T __attribute__((ext_vector_type(N))) v) {
+  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 __attribute__((ext_vector_type(N / 2))) reduced = v.hi + v.lo;
+      T [[clang::ext_vector_type(N / 2)]] reduced = v.hi + v.lo;
       return simd_reduce(reduced);
     }
   }
@@ -1144,6 +1144,8 @@ accepted values.
   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.
   }];
 }
 

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 __attribute__((ext_vector_type(2))) vec2_a, vec2_b;
+  __bf16 [[clang::ext_vector_type(2)]] vec2_a, vec2_b;
   vec2_a = vec2_b;
 
-  __bf16 __attribute__((ext_vector_type(4))) vec4_a, vec4_b;
+  __bf16 [[clang::ext_vector_type(4)]] vec4_a, vec4_b;
   vec4_a = vec4_b;
 
-  __bf16 __attribute__((ext_vector_type(8))) vec8_a, vec8_b;
+  __bf16 [[clang::ext_vector_type(8)]] vec8_a, vec8_b;
   vec8_a = vec8_b;
 
-  __bf16 __attribute__((ext_vector_type(16))) vec16_a, vec16_b;
+  __bf16 [[clang::ext_vector_type(16)]] vec16_a, vec16_b;
   vec16_a = vec16_b;
 }