-
Notifications
You must be signed in to change notification settings - Fork 10.8k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
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
- Loading branch information
Showing
8 changed files
with
199 additions
and
48 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -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 | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -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); | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.