[clang] [WIP][Clang] Allow floating point fixed vectors with atomic builtins (PR #129495)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 21 04:24:59 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
@llvm/pr-subscribers-clang-codegen
Author: Vikram Hegde (vikramRH)
<details>
<summary>Changes</summary>
https://github.com/llvm/llvm-project/pull/86796 added support for atomicrmw FP ops with fixed vector types. This patch intends to allow the same with clang atomic builtins. Any comments/concerns here would be helpful..
---
Full diff: https://github.com/llvm/llvm-project/pull/129495.diff
7 Files Affected:
- (modified) clang/include/clang/AST/Type.h (+3)
- (modified) clang/lib/AST/Type.cpp (+11)
- (modified) clang/lib/CodeGen/CGAtomic.cpp (+14-20)
- (modified) clang/lib/Sema/SemaChecking.cpp (+6-5)
- (modified) clang/test/CodeGen/fp-atomic-ops.c (+22)
- (modified) clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu (+51)
- (modified) clang/test/Sema/atomic-ops.c (+36-8)
``````````diff
diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h
index c3ff7ebd88516..34f0037e83efc 100644
--- a/clang/include/clang/AST/Type.h
+++ b/clang/include/clang/AST/Type.h
@@ -2738,6 +2738,9 @@ class alignas(TypeAlignment) Type : public ExtQualsTypeCommonBase {
/// Determine wither this type is a C++ elaborated-type-specifier.
bool isElaboratedTypeSpecifier() const;
+ // check whether the type is compatible with fp atomics.
+ bool isFPAtomicCompatibleType() const;
+
bool canDecayToPointerType() const;
/// Whether this type is represented natively as a pointer. This includes
diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp
index 8c11ec2e1fe24..3b082443d0ce3 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -2312,6 +2312,17 @@ bool Type::isRealType() const {
return isBitIntType();
}
+bool Type::isFPAtomicCompatibleType() const {
+ if (isa<ComplexType>(CanonicalType))
+ return false;
+ if (const auto *CVT = dyn_cast<VectorType>(CanonicalType)) {
+ if (CVT->isSizelessVectorType())
+ return false;
+ return CVT->getElementType()->isFPAtomicCompatibleType();
+ }
+ return isFloatingType();
+}
+
bool Type::isArithmeticType() const {
if (const auto *BT = dyn_cast<BuiltinType>(CanonicalType))
return BT->getKind() >= BuiltinType::Bool &&
diff --git a/clang/lib/CodeGen/CGAtomic.cpp b/clang/lib/CodeGen/CGAtomic.cpp
index 3adb2a7ad207f..776e989ef46cd 100644
--- a/clang/lib/CodeGen/CGAtomic.cpp
+++ b/clang/lib/CodeGen/CGAtomic.cpp
@@ -531,6 +531,7 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
bool PostOpMinMax = false;
unsigned PostOp = 0;
+ bool IsFloat = E->getValueType()->isFPAtomicCompatibleType();
switch (E->getOp()) {
case AtomicExpr::AO__c11_atomic_init:
case AtomicExpr::AO__opencl_atomic_init:
@@ -620,30 +621,26 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
case AtomicExpr::AO__atomic_add_fetch:
case AtomicExpr::AO__scoped_atomic_add_fetch:
- PostOp = E->getValueType()->isFloatingType() ? llvm::Instruction::FAdd
- : llvm::Instruction::Add;
+ PostOp = IsFloat ? llvm::Instruction::FAdd : llvm::Instruction::Add;
[[fallthrough]];
case AtomicExpr::AO__c11_atomic_fetch_add:
case AtomicExpr::AO__hip_atomic_fetch_add:
case AtomicExpr::AO__opencl_atomic_fetch_add:
case AtomicExpr::AO__atomic_fetch_add:
case AtomicExpr::AO__scoped_atomic_fetch_add:
- Op = E->getValueType()->isFloatingType() ? llvm::AtomicRMWInst::FAdd
- : llvm::AtomicRMWInst::Add;
+ Op = IsFloat ? llvm::AtomicRMWInst::FAdd : llvm::AtomicRMWInst::Add;
break;
case AtomicExpr::AO__atomic_sub_fetch:
case AtomicExpr::AO__scoped_atomic_sub_fetch:
- PostOp = E->getValueType()->isFloatingType() ? llvm::Instruction::FSub
- : llvm::Instruction::Sub;
+ PostOp = IsFloat ? llvm::Instruction::FSub : llvm::Instruction::Sub;
[[fallthrough]];
case AtomicExpr::AO__c11_atomic_fetch_sub:
case AtomicExpr::AO__hip_atomic_fetch_sub:
case AtomicExpr::AO__opencl_atomic_fetch_sub:
case AtomicExpr::AO__atomic_fetch_sub:
case AtomicExpr::AO__scoped_atomic_fetch_sub:
- Op = E->getValueType()->isFloatingType() ? llvm::AtomicRMWInst::FSub
- : llvm::AtomicRMWInst::Sub;
+ Op = IsFloat ? llvm::AtomicRMWInst::FSub : llvm::AtomicRMWInst::Sub;
break;
case AtomicExpr::AO__atomic_min_fetch:
@@ -655,11 +652,10 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
case AtomicExpr::AO__opencl_atomic_fetch_min:
case AtomicExpr::AO__atomic_fetch_min:
case AtomicExpr::AO__scoped_atomic_fetch_min:
- Op = E->getValueType()->isFloatingType()
- ? llvm::AtomicRMWInst::FMin
- : (E->getValueType()->isSignedIntegerType()
- ? llvm::AtomicRMWInst::Min
- : llvm::AtomicRMWInst::UMin);
+ Op = IsFloat ? llvm::AtomicRMWInst::FMin
+ : (E->getValueType()->isSignedIntegerType()
+ ? llvm::AtomicRMWInst::Min
+ : llvm::AtomicRMWInst::UMin);
break;
case AtomicExpr::AO__atomic_max_fetch:
@@ -671,11 +667,10 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
case AtomicExpr::AO__opencl_atomic_fetch_max:
case AtomicExpr::AO__atomic_fetch_max:
case AtomicExpr::AO__scoped_atomic_fetch_max:
- Op = E->getValueType()->isFloatingType()
- ? llvm::AtomicRMWInst::FMax
- : (E->getValueType()->isSignedIntegerType()
- ? llvm::AtomicRMWInst::Max
- : llvm::AtomicRMWInst::UMax);
+ Op = IsFloat ? llvm::AtomicRMWInst::FMax
+ : (E->getValueType()->isSignedIntegerType()
+ ? llvm::AtomicRMWInst::Max
+ : llvm::AtomicRMWInst::UMax);
break;
case AtomicExpr::AO__atomic_and_fetch:
@@ -984,9 +979,8 @@ RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E) {
case AtomicExpr::AO__scoped_atomic_max_fetch:
case AtomicExpr::AO__scoped_atomic_min_fetch:
case AtomicExpr::AO__scoped_atomic_sub_fetch:
- ShouldCastToIntPtrTy = !MemTy->isFloatingType();
+ ShouldCastToIntPtrTy = !MemTy->isFPAtomicCompatibleType();
[[fallthrough]];
-
case AtomicExpr::AO__atomic_fetch_and:
case AtomicExpr::AO__atomic_fetch_nand:
case AtomicExpr::AO__atomic_fetch_or:
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index f9926c6b4adab..d1021cfef764e 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -3758,7 +3758,7 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
enum ArithOpExtraValueType {
AOEVT_None = 0,
AOEVT_Pointer = 1,
- AOEVT_FP = 2,
+ AOEVT_FPorFPVec = 2,
};
unsigned ArithAllows = AOEVT_None;
@@ -3804,7 +3804,7 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
case AtomicExpr::AO__opencl_atomic_fetch_sub:
case AtomicExpr::AO__hip_atomic_fetch_add:
case AtomicExpr::AO__hip_atomic_fetch_sub:
- ArithAllows = AOEVT_Pointer | AOEVT_FP;
+ ArithAllows = AOEVT_Pointer | AOEVT_FPorFPVec;
Form = Arithmetic;
break;
case AtomicExpr::AO__atomic_fetch_max:
@@ -3821,7 +3821,7 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
case AtomicExpr::AO__opencl_atomic_fetch_min:
case AtomicExpr::AO__hip_atomic_fetch_max:
case AtomicExpr::AO__hip_atomic_fetch_min:
- ArithAllows = AOEVT_FP;
+ ArithAllows = AOEVT_FPorFPVec;
Form = Arithmetic;
break;
case AtomicExpr::AO__c11_atomic_fetch_and:
@@ -3982,7 +3982,8 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
return true;
if (ValType->isPointerType())
return AllowedType & AOEVT_Pointer;
- if (!(ValType->isFloatingType() && (AllowedType & AOEVT_FP)))
+ if (!(ValType->isFPAtomicCompatibleType() &&
+ (AllowedType & AOEVT_FPorFPVec)))
return false;
// LLVM Parser does not allow atomicrmw with x86_fp80 type.
if (ValType->isSpecificBuiltinType(BuiltinType::LongDouble) &&
@@ -3992,7 +3993,7 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
return true;
};
if (!IsAllowedValueType(ValType, ArithAllows)) {
- auto DID = ArithAllows & AOEVT_FP
+ auto DID = ArithAllows & AOEVT_FPorFPVec
? (ArithAllows & AOEVT_Pointer
? diag::err_atomic_op_needs_atomic_int_ptr_or_fp
: diag::err_atomic_op_needs_atomic_int_or_fp)
diff --git a/clang/test/CodeGen/fp-atomic-ops.c b/clang/test/CodeGen/fp-atomic-ops.c
index c894e7b4ade37..0e17c3278fbee 100644
--- a/clang/test/CodeGen/fp-atomic-ops.c
+++ b/clang/test/CodeGen/fp-atomic-ops.c
@@ -27,6 +27,9 @@ typedef enum memory_order {
memory_order_seq_cst = __ATOMIC_SEQ_CST
} memory_order;
+typedef float float2 __attribute__((ext_vector_type(2)));
+typedef double double2 __attribute__((ext_vector_type(2)));
+
void test(float *f, float ff, double *d, double dd) {
// FLOAT: atomicrmw fadd ptr {{.*}} monotonic
__atomic_fetch_add(f, ff, memory_order_relaxed);
@@ -42,3 +45,22 @@ void test(float *f, float ff, double *d, double dd) {
__atomic_fetch_sub(d, dd, memory_order_relaxed);
#endif
}
+
+typedef float float2 __attribute__((ext_vector_type(2)));
+typedef double double2 __attribute__((ext_vector_type(2)));
+
+void test_vector(float2 *f, float2 ff, double2 *d, double2 dd) {
+ // FLOAT: atomicrmw fadd ptr {{.*}} <2 x float> {{.*}} monotonic
+ __atomic_fetch_add(f, ff, memory_order_relaxed);
+
+ // FLOAT: atomicrmw fsub ptr {{.*}} <2 x float> {{.*}} monotonic
+ __atomic_fetch_sub(f, ff, memory_order_relaxed);
+
+#ifdef DOUBLE
+ // DOUBLE: atomicrmw fadd ptr {{.*}} <2 x double> {{.*}} monotonic
+ __atomic_fetch_add(d, dd, memory_order_relaxed);
+
+ // DOUBLE: atomicrmw fsub ptr {{.*}} <2 x double> {{.*}} monotonic
+ __atomic_fetch_sub(d, dd, memory_order_relaxed);
+#endif
+}
diff --git a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
index 37fca614c3111..6afb39d6f8405 100644
--- a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -20,6 +20,8 @@
#include "Inputs/cuda.h"
#include <stdatomic.h>
+typedef float __attribute__((ext_vector_type(2))) vector_float;
+
__global__ void ffp1(float *p) {
// CHECK-LABEL: @_Z4ffp1Pf
// SAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4{{$}}
@@ -225,6 +227,55 @@ __global__ void ffp6(_Float16 *p) {
__hip_atomic_fetch_min(p, 1.0f, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
}
+__global__ void ffp7(vector_float *p) {
+ // CHECK-LABEL: @_Z4ffp7PDv2_f
+ // SAFEIR: atomicrmw fadd ptr {{.*}}<2 x float>{{.*}} monotonic, align 8{{$}}
+ // SAFEIR: atomicrmw fsub ptr {{.*}}<2 x float>{{.*}} monotonic, align 8{{$}}
+ // SAFEIR: atomicrmw fmax ptr {{.*}}<2 x float>{{.*}} monotonic, align 8{{$}}
+ // SAFEIR: atomicrmw fmin ptr {{.*}}<2 x float>{{.*}} monotonic, align 8{{$}}
+ // SAFEIR: atomicrmw fadd ptr {{.*}}<2 x float>{{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
+ // SAFEIR: atomicrmw fsub ptr {{.*}}<2 x float>{{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
+ // SAFEIR: atomicrmw fmax ptr {{.*}}<2 x float>{{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
+ // SAFEIR: atomicrmw fmin ptr {{.*}}<2 x float>{{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
+
+ // UNSAFEIR: atomicrmw fadd ptr {{.*}}<2 x float>{{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fsub ptr {{.*}}<2 x float>{{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fmax ptr {{.*}}<2 x float>{{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fmin ptr {{.*}}<2 x float>{{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fadd ptr {{.*}}<2 x float>{{.*}} monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fsub ptr {{.*}}<2 x float>{{.*}} monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fmax ptr {{.*}}<2 x float>{{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+ // UNSAFEIR: atomicrmw fmin ptr {{.*}}<2 x float>{{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
+
+ // SAFE: _Z4ffp7PDv2_f
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+ // SAFE: global_atomic_cmpswap
+
+ // UNSAFE: _Z4ffp7PDv2_f
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ // UNSAFE: global_atomic_cmpswap
+ __atomic_fetch_add(p, {1.0f, 1.0f}, memory_order_relaxed);
+ __atomic_fetch_sub(p, {1.0f, 1.0f}, memory_order_relaxed);
+ __atomic_fetch_max(p, {1.0f, 1.0f}, memory_order_relaxed);
+ __atomic_fetch_min(p, {1.0f, 1.0f}, memory_order_relaxed);
+ __hip_atomic_fetch_add(p, {1.0f, 1.0f}, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
+ __hip_atomic_fetch_sub(p, {1.0f, 1.0f}, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
+ __hip_atomic_fetch_max(p, {1.0f, 1.0f}, memory_order_relaxed, __HIP_MEMORY_SCOPE_AGENT);
+ __hip_atomic_fetch_min(p, {1.0f, 1.0f}, memory_order_relaxed, __HIP_MEMORY_SCOPE_WORKGROUP);
+}
+
// CHECK-LABEL: @_Z12test_cmpxchgPiii
// CHECK: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} acquire acquire, align 4{{$}}
// CHECK: cmpxchg weak ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} acquire acquire, align 4{{$}}
diff --git a/clang/test/Sema/atomic-ops.c b/clang/test/Sema/atomic-ops.c
index 725a12060d4e0..03d57517d1571 100644
--- a/clang/test/Sema/atomic-ops.c
+++ b/clang/test/Sema/atomic-ops.c
@@ -1,19 +1,19 @@
-// RUN: %clang_cc1 %s -verify=expected,fp80,noi128 -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify=expected,fp80,noi128 -fenable-matrix -fgnuc-version=4.2.1 -ffreestanding \
// RUN: -fsyntax-only -triple=i686-linux-gnu -std=c11
-// RUN: %clang_cc1 %s -verify=expected,noi128 -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify=expected,noi128 -fenable-matrix -fgnuc-version=4.2.1 -ffreestanding \
// RUN: -fsyntax-only -triple=i686-linux-android -std=c11
-// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -fenable-matrix -ffreestanding \
// RUN: -fsyntax-only -triple=powerpc64-linux-gnu -std=c11
-// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -fenable-matrix -ffreestanding \
// RUN: -fsyntax-only -triple=powerpc64-linux-gnu -std=c11 \
// RUN: -target-cpu pwr7
-// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -fenable-matrix -ffreestanding \
// RUN: -fsyntax-only -triple=powerpc64le-linux-gnu -std=c11 \
// RUN: -target-cpu pwr8 -DPPC64_PWR8
-// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -fenable-matrix -ffreestanding \
// RUN: -fsyntax-only -triple=powerpc64-unknown-aix -std=c11 \
// RUN: -target-cpu pwr8
-// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding \
+// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -fenable-matrix -ffreestanding \
// RUN: -fsyntax-only -triple=powerpc64-unknown-aix -std=c11 \
// RUN: -mabi=quadword-atomics -target-cpu pwr8 -DPPC64_PWR8
@@ -147,7 +147,15 @@ _Static_assert(__atomic_always_lock_free(2, (int[2]){}), "");
void dummyfn();
_Static_assert(__atomic_always_lock_free(2, dummyfn) || 1, "");
+typedef _Atomic(float __attribute__((vector_size(16)))) atomic_vector_float;
+typedef _Atomic(double __attribute__((vector_size(16)))) atomic_vector_double;
+typedef _Atomic(int __attribute__((vector_size(16)))) atomic_vector_int;
+typedef float __attribute__((ext_vector_type(4))) vector_float;
+typedef double __attribute__((ext_vector_type(2))) vector_double;
+typedef int __attribute__((ext_vector_type(4))) vector_int;
+typedef float float_mat_5x5 __attribute__((matrix_type(5, 5)));
+typedef _Complex double ComplexDouble;
#define _AS1 __attribute__((address_space(1)))
#define _AS2 __attribute__((address_space(2)))
@@ -156,7 +164,11 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
_Atomic(int*) *p, _Atomic(float) *f, _Atomic(double) *d,
_Atomic(long double) *ld,
int *I, const int *CI,
- int **P, float *F, double *D, struct S *s1, struct S *s2) {
+ int **P, float *F, double *D, struct S *s1, struct S *s2,
+ atomic_vector_float* vf, atomic_vector_double* vd,
+ atomic_vector_int* vi, vector_float* evf,
+ vector_double* evd, vector_int* evi, float_mat_5x5* fm,
+ ComplexDouble* cd) {
__c11_atomic_init(I, 5); // expected-error {{pointer to _Atomic}}
__c11_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const _Atomic(int) *' invalid)}}
@@ -224,6 +236,13 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
__c11_atomic_fetch_add(f, 1.0f, memory_order_seq_cst);
__c11_atomic_fetch_add(d, 1.0, memory_order_seq_cst);
__c11_atomic_fetch_add(ld, 1.0, memory_order_seq_cst); // fp80-error {{must be a pointer to atomic integer, pointer or supported floating point type}}
+
+ vector_float fvec = {1.0f, 1.0f, 1.0f, 1.0f};
+ vector_double dvec = {1.0, 1.0};
+ vector_int ivec = {1, 1, 1, 1};
+ __c11_atomic_fetch_add(vf, fvec, memory_order_seq_cst);
+ __c11_atomic_fetch_add(vd, dvec, memory_order_seq_cst);
+ __c11_atomic_fetch_add(vi, ivec, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer, pointer or supported floating point type}}
__c11_atomic_fetch_min(i, 1, memory_order_seq_cst);
__c11_atomic_fetch_min(p, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer or supported floating point type}}
__c11_atomic_fetch_min(f, 1.0f, memory_order_seq_cst);
@@ -240,6 +259,15 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
__atomic_fetch_sub(P, 3, memory_order_seq_cst);
__atomic_fetch_sub(F, 3, memory_order_seq_cst);
__atomic_fetch_sub(s1, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer, pointer or supported floating point type}}
+ __atomic_fetch_sub(evf, fvec, memory_order_seq_cst);
+ __atomic_fetch_sub(evd, dvec, memory_order_seq_cst);
+ __atomic_fetch_sub(evi, ivec, memory_order_seq_cst); // expected-error {{must be a pointer to integer, pointer or supported floating point type}}
+
+ float_mat_5x5 f1;
+ __atomic_fetch_sub(fm, f1, memory_order_seq_cst); // expected-error {{must be a pointer to integer, pointer or supported floating point type}}
+ ComplexDouble f2 = {1.0, 2.0};
+ __atomic_fetch_sub(cd, f2, memory_order_seq_cst); // expected-error {{must be a pointer to integer, pointer or supported floating point type}}
+
__atomic_fetch_min(F, 3, memory_order_seq_cst);
__atomic_fetch_min(D, 3, memory_order_seq_cst);
__atomic_fetch_max(F, 3, memory_order_seq_cst);
``````````
</details>
https://github.com/llvm/llvm-project/pull/129495
More information about the cfe-commits
mailing list