yaxunl updated this revision to Diff 283067.
yaxunl added a comment.
Herald added subscribers: atanasyan, sdardis.
added tests for targets supporting fp atomics.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D71726/new/
https://reviews.llvm.org/D71726
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/TargetInfo.h
clang/lib/Basic/Targets/AArch64.h
clang/lib/Basic/Targets/AMDGPU.h
clang/lib/Basic/Targets/ARM.h
clang/lib/Basic/Targets/Hexagon.h
clang/lib/Basic/Targets/Mips.h
clang/lib/Basic/Targets/X86.h
clang/lib/CodeGen/CGAtomic.cpp
clang/lib/Sema/SemaChecking.cpp
clang/test/CodeGen/fp-atomic-ops.c
clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
clang/test/CodeGenOpenCL/atomic-ops.cl
clang/test/Sema/atomic-ops.c
clang/test/SemaCUDA/amdgpu-atomic-ops.cu
clang/test/SemaOpenCL/atomic-ops.cl
Index: clang/test/SemaOpenCL/atomic-ops.cl
===================================================================
--- clang/test/SemaOpenCL/atomic-ops.cl
+++ clang/test/SemaOpenCL/atomic-ops.cl
@@ -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 @@
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 @@
__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.
Index: clang/test/SemaCUDA/amdgpu-atomic-ops.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/amdgpu-atomic-ops.cu
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 %s -verify -fsyntax-only -triple=amdgcn-amd-amdhsa \
+// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
+// RUN: -fnative-half-arguments-and-returns
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+__device__ _Float16 test_Flot16(_Float16 *p) {
+ return __atomic_fetch_sub(p, 1.0f16, memory_order_relaxed);
+ // expected-error@-1 {{address argument to atomic operation must be a pointer to integer, pointer or supported floating point type ('_Float16 *' invalid)}}
+}
+
+__device__ __fp16 test_fp16(__fp16 *p) {
+ return __atomic_fetch_sub(p, 1.0f16, memory_order_relaxed);
+ // expected-error@-1 {{address argument to atomic operation must be a pointer to integer, pointer or supported floating point type ('__fp16 *' invalid)}}
+}
Index: clang/test/Sema/atomic-ops.c
===================================================================
--- clang/test/Sema/atomic-ops.c
+++ clang/test/Sema/atomic-ops.c
@@ -99,7 +99,8 @@
#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 +115,7 @@
__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 +138,7 @@
__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 +167,22 @@
__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); // expected-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 +192,12 @@
_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);
Index: clang/test/CodeGenOpenCL/atomic-ops.cl
===================================================================
--- clang/test/CodeGenOpenCL/atomic-ops.cl
+++ clang/test/CodeGenOpenCL/atomic-ops.cl
@@ -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 @@
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()
{
Index: clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -0,0 +1,27 @@
+// 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);
+}
Index: clang/test/CodeGen/fp-atomic-ops.c
===================================================================
--- /dev/null
+++ clang/test/CodeGen/fp-atomic-ops.c
@@ -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
+}
Index: clang/lib/Sema/SemaChecking.cpp
===================================================================
--- clang/lib/Sema/SemaChecking.cpp
+++ clang/lib/Sema/SemaChecking.cpp
@@ -4703,9 +4703,11 @@
// 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)
+ if (IsAddSub && !ValType->isIntegerType() && !ValType->isPointerType() &&
+ (!ValType->isFloatingType() ||
+ !Context.getTargetInfo().isFPAtomicFetchAddSubSupported(
+ Context.getFloatTypeSemantics(ValType)))) {
+ Diag(ExprRange.getBegin(), diag::err_atomic_op_needs_atomic_int_ptr_or_fp)
<< IsC11 << Ptr->getType() << Ptr->getSourceRange();
return ExprError();
}
@@ -4832,7 +4834,8 @@
// 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 == Init || (Form == Arithmetic && ValType->isIntegerType()) ||
+ (IsAddSub && ValType->isFloatingType()))
Ty = ValType;
else if (Form == Copy || Form == Xchg) {
if (IsPassedByAddress) {
Index: clang/lib/CodeGen/CGAtomic.cpp
===================================================================
--- clang/lib/CodeGen/CGAtomic.cpp
+++ clang/lib/CodeGen/CGAtomic.cpp
@@ -595,21 +595,25 @@
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:
@@ -807,6 +811,7 @@
bool Oversized = getContext().toBits(sizeChars) > MaxInlineWidthInBits;
bool Misaligned = (Ptr.getAlignment() % sizeChars) != 0;
bool UseLibcall = Misaligned | Oversized;
+ bool ShouldCastToIntPtrTy = true;
if (UseLibcall) {
CGM.getDiags().Report(E->getBeginLoc(), diag::warn_atomic_op_misaligned)
@@ -876,11 +881,14 @@
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:
@@ -921,15 +929,23 @@
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) {
Index: clang/lib/Basic/Targets/X86.h
===================================================================
--- clang/lib/Basic/Targets/X86.h
+++ clang/lib/Basic/Targets/X86.h
@@ -347,6 +347,18 @@
uint64_t getPointerAlignV(unsigned AddrSpace) const override {
return getPointerWidthV(AddrSpace);
}
+
+ bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const override {
+ switch (llvm::APFloat::SemanticsToEnum(FS)) {
+ case llvm::APFloat::S_IEEEsingle:
+ return MaxAtomicInlineWidth >= 32;
+ case llvm::APFloat::S_IEEEdouble:
+ return MaxAtomicInlineWidth >= 64;
+ default:
+ return false;
+ }
+ }
};
// X86-32 generic target
Index: clang/lib/Basic/Targets/Mips.h
===================================================================
--- clang/lib/Basic/Targets/Mips.h
+++ clang/lib/Basic/Targets/Mips.h
@@ -407,6 +407,17 @@
bool validateTarget(DiagnosticsEngine &Diags) const override;
bool hasExtIntType() const override { return true; }
+
+ bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const override {
+ switch (llvm::APFloat::SemanticsToEnum(FS)) {
+ case llvm::APFloat::S_IEEEsingle:
+ case llvm::APFloat::S_IEEEdouble:
+ return true;
+ default:
+ return false;
+ }
+ }
};
} // namespace targets
} // namespace clang
Index: clang/lib/Basic/Targets/Hexagon.h
===================================================================
--- clang/lib/Basic/Targets/Hexagon.h
+++ clang/lib/Basic/Targets/Hexagon.h
@@ -140,6 +140,17 @@
}
bool hasExtIntType() const override { return true; }
+
+ bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const override {
+ switch (llvm::APFloat::SemanticsToEnum(FS)) {
+ case llvm::APFloat::S_IEEEsingle:
+ case llvm::APFloat::S_IEEEdouble:
+ return true;
+ default:
+ return false;
+ }
+ }
};
} // namespace targets
} // namespace clang
Index: clang/lib/Basic/Targets/ARM.h
===================================================================
--- clang/lib/Basic/Targets/ARM.h
+++ clang/lib/Basic/Targets/ARM.h
@@ -188,6 +188,16 @@
bool hasExtIntType() const override { return true; }
const char *getBFloat16Mangling() const override { return "u6__bf16"; };
+
+ bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const override {
+ switch (llvm::APFloat::SemanticsToEnum(FS)) {
+ case llvm::APFloat::S_IEEEsingle:
+ return true;
+ default:
+ return false;
+ }
+ }
};
class LLVM_LIBRARY_VISIBILITY ARMleTargetInfo : public ARMTargetInfo {
Index: clang/lib/Basic/Targets/AMDGPU.h
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.h
+++ clang/lib/Basic/Targets/AMDGPU.h
@@ -389,6 +389,17 @@
void setAuxTarget(const TargetInfo *Aux) override;
bool hasExtIntType() const override { return true; }
+
+ bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const override {
+ switch (llvm::APFloat::SemanticsToEnum(FS)) {
+ case llvm::APFloat::S_IEEEsingle:
+ case llvm::APFloat::S_IEEEdouble:
+ return true;
+ default:
+ return false;
+ }
+ }
};
} // namespace targets
Index: clang/lib/Basic/Targets/AArch64.h
===================================================================
--- clang/lib/Basic/Targets/AArch64.h
+++ clang/lib/Basic/Targets/AArch64.h
@@ -130,6 +130,17 @@
bool hasInt128Type() const override;
bool hasExtIntType() const override { return true; }
+
+ bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const override {
+ switch (llvm::APFloat::SemanticsToEnum(FS)) {
+ case llvm::APFloat::S_IEEEsingle:
+ case llvm::APFloat::S_IEEEdouble:
+ return true;
+ default:
+ return false;
+ }
+ }
};
class LLVM_LIBRARY_VISIBILITY AArch64leTargetInfo : public AArch64TargetInfo {
Index: clang/include/clang/Basic/TargetInfo.h
===================================================================
--- clang/include/clang/Basic/TargetInfo.h
+++ clang/include/clang/Basic/TargetInfo.h
@@ -1455,6 +1455,12 @@
/// Whether target allows debuginfo types for decl only variables.
virtual bool allowDebugInfoForExternalVar() const { return false; }
+ /// Whether floating point atomic fetch add/sub is supported.
+ virtual bool
+ isFPAtomicFetchAddSubSupported(const llvm::fltSemantics &FS) const {
+ return false;
+ }
+
protected:
/// Copy type and layout related info.
void copyAuxTarget(const TargetInfo *Aux);
Index: clang/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -7961,6 +7961,9 @@
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)">;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
