forked from OSchip/llvm-project
Let clang atomic builtins fetch add/sub support floating point types
Recently atomicrmw started to support fadd/fsub: https://reviews.llvm.org/D53965 However clang atomic builtins fetch add/sub still does not support emitting atomicrmw fadd/fsub. This patch adds that. Reviewed by: John McCall, Artem Belevich, Matt Arsenault, JF Bastien, James Y Knight, Louis Dionne, Olivier Giroux Differential Revision: https://reviews.llvm.org/D71726
This commit is contained in:
Yaxun (Sam) Liu
parent
ddebed8e97
commit
61d065e21f
|
|
@ -8188,6 +8188,9 @@ def err_atomic_op_needs_non_const_pointer : Error<
|
|||
def err_atomic_op_needs_trivial_copy : Error<
|
||||
"address argument to atomic operation must be a pointer to a "
|
||||
"trivially-copyable type (%0 invalid)">;
|
||||
def err_atomic_op_needs_atomic_int_ptr_or_fp : Error<
|
||||
"address argument to atomic operation must be a pointer to %select{|atomic }0"
|
||||
"integer, pointer or supported floating point type (%1 invalid)">;
|
||||
def err_atomic_op_needs_atomic_int_or_ptr : Error<
|
||||
"address argument to atomic operation must be a pointer to %select{|atomic }0"
|
||||
"integer or pointer (%1 invalid)">;
|
||||
|
|
|
|||
|
|
@ -602,21 +602,25 @@ static void EmitAtomicOp(CodeGenFunction &CGF, AtomicExpr *E, Address Dest,
|
|||
break;
|
||||
|
||||
case AtomicExpr::AO__atomic_add_fetch:
|
||||
PostOp = llvm::Instruction::Add;
|
||||
PostOp = E->getValueType()->isFloatingType() ? llvm::Instruction::FAdd
|
||||
: llvm::Instruction::Add;
|
||||
LLVM_FALLTHROUGH;
|
||||
case AtomicExpr::AO__c11_atomic_fetch_add:
|
||||
case AtomicExpr::AO__opencl_atomic_fetch_add:
|
||||
case AtomicExpr::AO__atomic_fetch_add:
|
||||
Op = llvm::AtomicRMWInst::Add;
|
||||
Op = E->getValueType()->isFloatingType() ? llvm::AtomicRMWInst::FAdd
|
||||
: llvm::AtomicRMWInst::Add;
|
||||
break;
|
||||
|
||||
case AtomicExpr::AO__atomic_sub_fetch:
|
||||
PostOp = llvm::Instruction::Sub;
|
||||
PostOp = E->getValueType()->isFloatingType() ? llvm::Instruction::FSub
|
||||
: llvm::Instruction::Sub;
|
||||
LLVM_FALLTHROUGH;
|
||||
case AtomicExpr::AO__c11_atomic_fetch_sub:
|
||||
case AtomicExpr::AO__opencl_atomic_fetch_sub:
|
||||
case AtomicExpr::AO__atomic_fetch_sub:
|
||||
Op = llvm::AtomicRMWInst::Sub;
|
||||
Op = E->getValueType()->isFloatingType() ? llvm::AtomicRMWInst::FSub
|
||||
: llvm::AtomicRMWInst::Sub;
|
||||
break;
|
||||
|
||||
case AtomicExpr::AO__atomic_min_fetch:
|
||||
|
|
@ -813,6 +817,8 @@ RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E) {
|
|||
bool Oversized = getContext().toBits(TInfo.Width) > MaxInlineWidthInBits;
|
||||
bool Misaligned = (Ptr.getAlignment() % TInfo.Width) != 0;
|
||||
bool UseLibcall = Misaligned | Oversized;
|
||||
bool ShouldCastToIntPtrTy = true;
|
||||
|
||||
CharUnits MaxInlineWidth =
|
||||
getContext().toCharUnitsFromBits(MaxInlineWidthInBits);
|
||||
|
||||
|
|
@ -892,11 +898,14 @@ RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E) {
|
|||
EmitStoreOfScalar(Val1Scalar, MakeAddrLValue(Temp, Val1Ty));
|
||||
break;
|
||||
}
|
||||
LLVM_FALLTHROUGH;
|
||||
LLVM_FALLTHROUGH;
|
||||
case AtomicExpr::AO__atomic_fetch_add:
|
||||
case AtomicExpr::AO__atomic_fetch_sub:
|
||||
case AtomicExpr::AO__atomic_add_fetch:
|
||||
case AtomicExpr::AO__atomic_sub_fetch:
|
||||
ShouldCastToIntPtrTy = !MemTy->isFloatingType();
|
||||
LLVM_FALLTHROUGH;
|
||||
|
||||
case AtomicExpr::AO__c11_atomic_store:
|
||||
case AtomicExpr::AO__c11_atomic_exchange:
|
||||
case AtomicExpr::AO__opencl_atomic_store:
|
||||
|
|
@ -937,15 +946,23 @@ RValue CodeGenFunction::EmitAtomicExpr(AtomicExpr *E) {
|
|||
LValue AtomicVal = MakeAddrLValue(Ptr, AtomicTy);
|
||||
AtomicInfo Atomics(*this, AtomicVal);
|
||||
|
||||
Ptr = Atomics.emitCastToAtomicIntPointer(Ptr);
|
||||
if (Val1.isValid()) Val1 = Atomics.convertToAtomicIntPointer(Val1);
|
||||
if (Val2.isValid()) Val2 = Atomics.convertToAtomicIntPointer(Val2);
|
||||
if (Dest.isValid())
|
||||
Dest = Atomics.emitCastToAtomicIntPointer(Dest);
|
||||
else if (E->isCmpXChg())
|
||||
if (ShouldCastToIntPtrTy) {
|
||||
Ptr = Atomics.emitCastToAtomicIntPointer(Ptr);
|
||||
if (Val1.isValid())
|
||||
Val1 = Atomics.convertToAtomicIntPointer(Val1);
|
||||
if (Val2.isValid())
|
||||
Val2 = Atomics.convertToAtomicIntPointer(Val2);
|
||||
}
|
||||
if (Dest.isValid()) {
|
||||
if (ShouldCastToIntPtrTy)
|
||||
Dest = Atomics.emitCastToAtomicIntPointer(Dest);
|
||||
} else if (E->isCmpXChg())
|
||||
Dest = CreateMemTemp(RValTy, "cmpxchg.bool");
|
||||
else if (!RValTy->isVoidType())
|
||||
Dest = Atomics.emitCastToAtomicIntPointer(Atomics.CreateTempAlloca());
|
||||
else if (!RValTy->isVoidType()) {
|
||||
Dest = Atomics.CreateTempAlloca();
|
||||
if (ShouldCastToIntPtrTy)
|
||||
Dest = Atomics.emitCastToAtomicIntPointer(Dest);
|
||||
}
|
||||
|
||||
// Use a library call. See: http://gcc.gnu.org/wiki/Atomic/GCCMM/LIbrary .
|
||||
if (UseLibcall) {
|
||||
|
|
|
|||
|
|
@ -4931,7 +4931,8 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
|
|||
case AtomicExpr::AO__atomic_add_fetch:
|
||||
case AtomicExpr::AO__atomic_sub_fetch:
|
||||
IsAddSub = true;
|
||||
LLVM_FALLTHROUGH;
|
||||
Form = Arithmetic;
|
||||
break;
|
||||
case AtomicExpr::AO__c11_atomic_fetch_and:
|
||||
case AtomicExpr::AO__c11_atomic_fetch_or:
|
||||
case AtomicExpr::AO__c11_atomic_fetch_xor:
|
||||
|
|
@ -4946,6 +4947,8 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
|
|||
case AtomicExpr::AO__atomic_or_fetch:
|
||||
case AtomicExpr::AO__atomic_xor_fetch:
|
||||
case AtomicExpr::AO__atomic_nand_fetch:
|
||||
Form = Arithmetic;
|
||||
break;
|
||||
case AtomicExpr::AO__c11_atomic_fetch_min:
|
||||
case AtomicExpr::AO__c11_atomic_fetch_max:
|
||||
case AtomicExpr::AO__opencl_atomic_fetch_min:
|
||||
|
|
@ -5038,10 +5041,24 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
|
|||
|
||||
// For an arithmetic operation, the implied arithmetic must be well-formed.
|
||||
if (Form == Arithmetic) {
|
||||
// gcc does not enforce these rules for GNU atomics, but we do so for sanity.
|
||||
if (IsAddSub && !ValType->isIntegerType()
|
||||
&& !ValType->isPointerType()) {
|
||||
Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int_or_ptr)
|
||||
// gcc does not enforce these rules for GNU atomics, but we do so for
|
||||
// sanity.
|
||||
auto IsAllowedValueType = [&](QualType ValType) {
|
||||
if (ValType->isIntegerType())
|
||||
return true;
|
||||
if (ValType->isPointerType())
|
||||
return true;
|
||||
if (!ValType->isFloatingType())
|
||||
return false;
|
||||
// LLVM Parser does not allow atomicrmw with x86_fp80 type.
|
||||
if (ValType->isSpecificBuiltinType(BuiltinType::LongDouble) &&
|
||||
&Context.getTargetInfo().getLongDoubleFormat() ==
|
||||
&llvm::APFloat::x87DoubleExtended())
|
||||
return false;
|
||||
return true;
|
||||
};
|
||||
if (IsAddSub && !IsAllowedValueType(ValType)) {
|
||||
Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int_ptr_or_fp)
|
||||
<< IsC11 << Ptr->getType() << Ptr->getSourceRange();
|
||||
return ExprError();
|
||||
}
|
||||
|
|
@ -5168,7 +5185,9 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
|
|||
// passed by address. For the rest, GNU uses by-address and C11 uses
|
||||
// by-value.
|
||||
assert(Form != Load);
|
||||
if (Form == Init || (Form == Arithmetic && ValType->isIntegerType()))
|
||||
if (Form == Arithmetic && ValType->isPointerType())
|
||||
Ty = Context.getPointerDiffType();
|
||||
else if (Form == Init || Form == Arithmetic)
|
||||
Ty = ValType;
|
||||
else if (Form == Copy || Form == Xchg) {
|
||||
if (IsPassedByAddress) {
|
||||
|
|
@ -5177,9 +5196,7 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
|
|||
ExprRange.getBegin());
|
||||
}
|
||||
Ty = ByValType;
|
||||
} else if (Form == Arithmetic)
|
||||
Ty = Context.getPointerDiffType();
|
||||
else {
|
||||
} else {
|
||||
Expr *ValArg = APIOrderedArgs[i];
|
||||
// The value pointer is always dereferenced, a nullptr is undefined.
|
||||
CheckNonNullArgument(*this, ValArg, ExprRange.getBegin());
|
||||
|
|
|
|||
|
|
@ -0,0 +1,44 @@
|
|||
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=amdgcn-amd-amdhsa \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
|
||||
|
||||
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=aarch64-linux-gnu \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
|
||||
|
||||
// RUN: %clang_cc1 %s -emit-llvm -O0 -o - -triple=armv8-apple-ios7.0 \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT %s
|
||||
|
||||
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=hexagon \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
|
||||
|
||||
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=mips64-mti-linux-gnu \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
|
||||
|
||||
// RUN: %clang_cc1 %s -emit-llvm -O0 -o - -triple=i686-linux-gnu \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT %s
|
||||
|
||||
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=x86_64-linux-gnu \
|
||||
// RUN: | opt -instnamer -S | FileCheck -check-prefixes=FLOAT,DOUBLE %s
|
||||
|
||||
typedef enum memory_order {
|
||||
memory_order_relaxed = __ATOMIC_RELAXED,
|
||||
memory_order_acquire = __ATOMIC_ACQUIRE,
|
||||
memory_order_release = __ATOMIC_RELEASE,
|
||||
memory_order_acq_rel = __ATOMIC_ACQ_REL,
|
||||
memory_order_seq_cst = __ATOMIC_SEQ_CST
|
||||
} memory_order;
|
||||
|
||||
void test(float *f, float ff, double *d, double dd) {
|
||||
// FLOAT: atomicrmw fadd float* {{.*}} monotonic
|
||||
__atomic_fetch_add(f, ff, memory_order_relaxed);
|
||||
|
||||
// FLOAT: atomicrmw fsub float* {{.*}} monotonic
|
||||
__atomic_fetch_sub(f, ff, memory_order_relaxed);
|
||||
|
||||
#ifdef DOUBLE
|
||||
// DOUBLE: atomicrmw fadd double* {{.*}} monotonic
|
||||
__atomic_fetch_add(d, dd, memory_order_relaxed);
|
||||
|
||||
// DOUBLE: atomicrmw fsub double* {{.*}} monotonic
|
||||
__atomic_fetch_sub(d, dd, memory_order_relaxed);
|
||||
#endif
|
||||
}
|
||||
|
|
@ -0,0 +1,41 @@
|
|||
// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
|
||||
// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
|
||||
// RUN: -fnative-half-arguments-and-returns | FileCheck %s
|
||||
|
||||
// REQUIRES: amdgpu-registered-target
|
||||
|
||||
#include "Inputs/cuda.h"
|
||||
#include <stdatomic.h>
|
||||
|
||||
__device__ float ffp1(float *p) {
|
||||
// CHECK-LABEL: @_Z4ffp1Pf
|
||||
// CHECK: atomicrmw fadd float* {{.*}} monotonic
|
||||
return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
|
||||
}
|
||||
|
||||
__device__ double ffp2(double *p) {
|
||||
// CHECK-LABEL: @_Z4ffp2Pd
|
||||
// CHECK: atomicrmw fsub double* {{.*}} monotonic
|
||||
return __atomic_fetch_sub(p, 1.0, memory_order_relaxed);
|
||||
}
|
||||
|
||||
// long double is the same as double for amdgcn.
|
||||
__device__ long double ffp3(long double *p) {
|
||||
// CHECK-LABEL: @_Z4ffp3Pe
|
||||
// CHECK: atomicrmw fsub double* {{.*}} monotonic
|
||||
return __atomic_fetch_sub(p, 1.0L, memory_order_relaxed);
|
||||
}
|
||||
|
||||
__device__ double ffp4(double *p, float f) {
|
||||
// CHECK-LABEL: @_Z4ffp4Pdf
|
||||
// CHECK: fpext float {{.*}} to double
|
||||
// CHECK: atomicrmw fsub double* {{.*}} monotonic
|
||||
return __atomic_fetch_sub(p, f, memory_order_relaxed);
|
||||
}
|
||||
|
||||
__device__ double ffp5(double *p, int i) {
|
||||
// CHECK-LABEL: @_Z4ffp5Pdi
|
||||
// CHECK: sitofp i32 {{.*}} to double
|
||||
// CHECK: atomicrmw fsub double* {{.*}} monotonic
|
||||
return __atomic_fetch_sub(p, i, memory_order_relaxed);
|
||||
}
|
||||
|
|
@ -1,12 +1,17 @@
|
|||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -O0 -o - -triple=amdgcn-amd-amdhsa-amdgizcl | opt -instnamer -S | FileCheck %s
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -O0 -o - -triple=amdgcn-amd-amdhsa \
|
||||
// RUN: | opt -instnamer -S | FileCheck %s
|
||||
|
||||
// Also test serialization of atomic operations here, to avoid duplicating the test.
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-pch -O0 -o %t -triple=amdgcn-amd-amdhsa-amdgizcl
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -include-pch %t -O0 -triple=amdgcn-amd-amdhsa-amdgizcl -emit-llvm -o - | opt -instnamer -S | FileCheck %s
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-pch -O0 -o %t -triple=amdgcn-amd-amdhsa
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -include-pch %t -O0 -triple=amdgcn-amd-amdhsa \
|
||||
// RUN: -emit-llvm -o - | opt -instnamer -S | FileCheck %s
|
||||
|
||||
#ifndef ALREADY_INCLUDED
|
||||
#define ALREADY_INCLUDED
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
|
||||
|
||||
typedef __INTPTR_TYPE__ intptr_t;
|
||||
typedef int int8 __attribute__((ext_vector_type(8)));
|
||||
|
||||
|
|
@ -185,6 +190,18 @@ float ff3(atomic_float *d) {
|
|||
return __opencl_atomic_exchange(d, 2, memory_order_seq_cst, memory_scope_work_group);
|
||||
}
|
||||
|
||||
float ff4(global atomic_float *d, float a) {
|
||||
// CHECK-LABEL: @ff4
|
||||
// CHECK: atomicrmw fadd float addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
|
||||
return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
|
||||
}
|
||||
|
||||
float ff5(global atomic_double *d, double a) {
|
||||
// CHECK-LABEL: @ff5
|
||||
// CHECK: atomicrmw fadd double addrspace(1)* {{.*}} syncscope("workgroup-one-as") monotonic
|
||||
return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
|
||||
}
|
||||
|
||||
// CHECK-LABEL: @atomic_init_foo
|
||||
void atomic_init_foo()
|
||||
{
|
||||
|
|
|
|||
|
|
@ -1,4 +1,9 @@
|
|||
// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding -fsyntax-only -triple=i686-linux-gnu -std=c11
|
||||
// RUN: %clang_cc1 %s -verify=expected,fp80,noi128 -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: -fsyntax-only -triple=i686-linux-android -std=c11
|
||||
// RUN: %clang_cc1 %s -verify -fgnuc-version=4.2.1 -ffreestanding \
|
||||
// RUN: -fsyntax-only -triple=powerpc64-linux-gnu -std=c11
|
||||
|
||||
// Basic parsing/Sema tests for __c11_atomic_*
|
||||
|
||||
|
|
@ -51,7 +56,7 @@ _Static_assert(atomic_is_lock_free((atomic_char*)0), "");
|
|||
_Static_assert(atomic_is_lock_free((atomic_short*)0), "");
|
||||
_Static_assert(atomic_is_lock_free((atomic_int*)0), "");
|
||||
_Static_assert(atomic_is_lock_free((atomic_long*)0), "");
|
||||
// expected-error@+1 {{__int128 is not supported on this target}}
|
||||
// noi128-error@+1 {{__int128 is not supported on this target}}
|
||||
_Static_assert(atomic_is_lock_free((_Atomic(__int128)*)0), ""); // expected-error {{not an integral constant expression}}
|
||||
_Static_assert(atomic_is_lock_free(0 + (atomic_char*)0), "");
|
||||
|
||||
|
|
@ -99,7 +104,8 @@ _Static_assert(__atomic_always_lock_free(8, &i64), "");
|
|||
#define _AS2 __attribute__((address_space(2)))
|
||||
|
||||
void f(_Atomic(int) *i, const _Atomic(int) *ci,
|
||||
_Atomic(int*) *p, _Atomic(float) *d,
|
||||
_Atomic(int*) *p, _Atomic(float) *f, _Atomic(double) *d,
|
||||
_Atomic(long double) *ld,
|
||||
int *I, const int *CI,
|
||||
int **P, float *D, struct S *s1, struct S *s2) {
|
||||
__c11_atomic_init(I, 5); // expected-error {{pointer to _Atomic}}
|
||||
|
|
@ -114,7 +120,7 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
|
|||
|
||||
__c11_atomic_load(i, memory_order_seq_cst);
|
||||
__c11_atomic_load(p, memory_order_seq_cst);
|
||||
__c11_atomic_load(d, memory_order_seq_cst);
|
||||
__c11_atomic_load(f, memory_order_seq_cst);
|
||||
__c11_atomic_load(ci, memory_order_seq_cst);
|
||||
|
||||
int load_n_1 = __atomic_load_n(I, memory_order_relaxed);
|
||||
|
|
@ -137,7 +143,7 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
|
|||
|
||||
__c11_atomic_store(i, 1, memory_order_seq_cst);
|
||||
__c11_atomic_store(p, 1, memory_order_seq_cst); // expected-warning {{incompatible integer to pointer conversion}}
|
||||
(int)__c11_atomic_store(d, 1, memory_order_seq_cst); // expected-error {{operand of type 'void'}}
|
||||
(int)__c11_atomic_store(f, 1, memory_order_seq_cst); // expected-error {{operand of type 'void'}}
|
||||
|
||||
__atomic_store_n(I, 4, memory_order_release);
|
||||
__atomic_store_n(I, 4.0, memory_order_release);
|
||||
|
|
@ -166,20 +172,22 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
|
|||
|
||||
__c11_atomic_fetch_add(i, 1, memory_order_seq_cst);
|
||||
__c11_atomic_fetch_add(p, 1, memory_order_seq_cst);
|
||||
__c11_atomic_fetch_add(d, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer or pointer}}
|
||||
__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}}
|
||||
|
||||
__atomic_fetch_add(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer or pointer}}
|
||||
__atomic_fetch_add(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer, pointer or supported floating point type}}
|
||||
__atomic_fetch_sub(I, 3, memory_order_seq_cst);
|
||||
__atomic_fetch_sub(P, 3, memory_order_seq_cst);
|
||||
__atomic_fetch_sub(D, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer or pointer}}
|
||||
__atomic_fetch_sub(s1, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer or pointer}}
|
||||
__atomic_fetch_sub(D, 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_min(D, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
|
||||
__atomic_fetch_max(P, 3, memory_order_seq_cst); // expected-error {{must be a pointer to integer}}
|
||||
__atomic_fetch_max(p, 3); // expected-error {{too few arguments to function call, expected 3, have 2}}
|
||||
|
||||
__c11_atomic_fetch_and(i, 1, memory_order_seq_cst);
|
||||
__c11_atomic_fetch_and(p, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer}}
|
||||
__c11_atomic_fetch_and(d, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer}}
|
||||
__c11_atomic_fetch_and(f, 1, memory_order_seq_cst); // expected-error {{must be a pointer to atomic integer}}
|
||||
|
||||
__atomic_fetch_and(i, 3, memory_order_seq_cst); // expected-error {{pointer to integer}}
|
||||
__atomic_fetch_or(I, 3, memory_order_seq_cst);
|
||||
|
|
@ -189,12 +197,12 @@ void f(_Atomic(int) *i, const _Atomic(int) *ci,
|
|||
|
||||
_Bool cmpexch_1 = __c11_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst);
|
||||
_Bool cmpexch_2 = __c11_atomic_compare_exchange_strong(p, P, (int*)1, memory_order_seq_cst, memory_order_seq_cst);
|
||||
_Bool cmpexch_3 = __c11_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
|
||||
_Bool cmpexch_3 = __c11_atomic_compare_exchange_strong(f, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
|
||||
(void)__c11_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
|
||||
|
||||
_Bool cmpexchw_1 = __c11_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst);
|
||||
_Bool cmpexchw_2 = __c11_atomic_compare_exchange_weak(p, P, (int*)1, memory_order_seq_cst, memory_order_seq_cst);
|
||||
_Bool cmpexchw_3 = __c11_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
|
||||
_Bool cmpexchw_3 = __c11_atomic_compare_exchange_weak(f, I, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{incompatible pointer types}}
|
||||
(void)__c11_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst); // expected-warning {{passing 'const int *' to parameter of type 'int *' discards qualifiers}}
|
||||
|
||||
_Bool cmpexch_4 = __atomic_compare_exchange_n(I, I, 5, 1, memory_order_seq_cst, memory_order_seq_cst);
|
||||
|
|
|
|||
|
|
@ -1,10 +1,13 @@
|
|||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only -triple=spir64
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only -triple=amdgcn-amdhsa-amd-opencl
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify=expected,spir \
|
||||
// RUN: -fsyntax-only -triple=spir64
|
||||
// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -fsyntax-only \
|
||||
// RUN: -triple=amdgcn-amd-amdhsa
|
||||
|
||||
// Basic parsing/Sema tests for __opencl_atomic_*
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
typedef __INTPTR_TYPE__ intptr_t;
|
||||
typedef int int8 __attribute__((ext_vector_type(8)));
|
||||
|
|
@ -36,7 +39,7 @@ int8 i64;
|
|||
|
||||
atomic_int gn;
|
||||
void f(atomic_int *i, const atomic_int *ci,
|
||||
atomic_intptr_t *p, atomic_float *d,
|
||||
atomic_intptr_t *p, atomic_float *f, atomic_double *d, atomic_half *h, // expected-error {{unknown type name 'atomic_half'}}
|
||||
int *I, const int *CI,
|
||||
intptr_t *P, float *D, struct S *s1, struct S *s2,
|
||||
global atomic_int *i_g, local atomic_int *i_l, private atomic_int *i_p,
|
||||
|
|
@ -57,37 +60,38 @@ void f(atomic_int *i, const atomic_int *ci,
|
|||
|
||||
__opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_load(p, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_load(d, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_load(f, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_load(i_c, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-constant _Atomic type ('__constant atomic_int *' (aka '__constant _Atomic(int) *') invalid)}}
|
||||
|
||||
__opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_store(p, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
(int)__opencl_atomic_store(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{operand of type 'void' where arithmetic or pointer type is required}}
|
||||
(int)__opencl_atomic_store(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{operand of type 'void' where arithmetic or pointer type is required}}
|
||||
|
||||
int exchange_1 = __opencl_atomic_exchange(i, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
int exchange_2 = __opencl_atomic_exchange(I, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to _Atomic}}
|
||||
|
||||
__opencl_atomic_fetch_add(i, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_add(p, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
__opencl_atomic_fetch_add(f, 1.0f, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_add(d, 1.0, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_and(p, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
__opencl_atomic_fetch_and(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
|
||||
__opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
|
||||
__opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
__opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
__opencl_atomic_fetch_min(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
__opencl_atomic_fetch_max(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
|
||||
|
||||
bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
|
||||
bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
|
||||
bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
|
||||
bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
|
||||
(void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}}
|
||||
|
||||
bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
|
||||
bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
|
||||
bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
|
||||
bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
|
||||
(void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}}
|
||||
|
||||
// Pointers to different address spaces are allowed.
|
||||
|
|
|
|||
Loading…
Reference in New Issue