llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Joseph Huber (jhuber6) <details> <summary>Changes</summary> 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 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; `````````` </details> https://github.com/llvm/llvm-project/pull/130177 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
