[clang] [llvm] [IR] Add fast-math flags to atomicrmw (PR #211267)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 22 08:05:28 PDT 2026
https://github.com/Lukacma updated https://github.com/llvm/llvm-project/pull/211267
>From 43aa1d1ba2bce4b932d9df9d08a7cf9f3190ee93 Mon Sep 17 00:00:00 2001
From: Marian Lukac <Marian.Lukac at arm.com>
Date: Wed, 22 Jul 2026 13:07:15 +0000
Subject: [PATCH] [IR] Add fast-math flags to atomicrmw
This patch optional fast-math flags to atomicrmw node.
The main reasoning for this is that there is floating
point arithmetic happening in this node so it makes
sense and we want to use these flags in AArch64
backend to enable fp unsafe codegen.
Note: This patch doesn't add support for `#pragma clang fp`
overrides.
---
clang/test/CodeGen/builtins-nvptx-ptx50.cu | 2 +-
clang/test/CodeGen/fp-atomic-ops.c | 5 +
clang/test/CodeGenCUDA/atomic-options.hip | 90 ++++-----
clang/test/CodeGenCUDA/builtins-amdgcn.cu | 8 +-
.../test/CodeGenCUDA/builtins-spirv-amdgcn.cu | 16 +-
.../builtins-unsafe-atomics-gfx90a.cu | 2 +-
...tins-unsafe-atomics-spirv-amdgcn-gfx90a.cu | 2 +-
clang/test/SemaHIP/incorrect-atomic-scope.hip | 2 +-
llvm/docs/LangRef.md | 22 ++-
llvm/include/llvm/Bitcode/LLVMBitCodes.h | 5 +-
.../CodeGen/GlobalISel/MachineIRBuilder.h | 4 +-
llvm/include/llvm/CodeGen/SelectionDAG.h | 6 +-
llvm/include/llvm/IR/IRBuilder.h | 27 ++-
llvm/include/llvm/IR/Instructions.h | 2 +-
llvm/include/llvm/IR/Operator.h | 5 +-
.../llvm/Transforms/Utils/LowerAtomic.h | 2 +-
llvm/lib/AsmParser/LLParser.cpp | 10 +-
llvm/lib/Bitcode/Reader/BitcodeReader.cpp | 44 +++--
llvm/lib/Bitcode/Writer/BitcodeWriter.cpp | 180 +++++++++---------
llvm/lib/CodeGen/AtomicExpandPass.cpp | 18 +-
llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp | 3 +-
.../CodeGen/GlobalISel/MachineIRBuilder.cpp | 4 +-
.../lib/CodeGen/SelectionDAG/SelectionDAG.cpp | 20 +-
.../SelectionDAG/SelectionDAGBuilder.cpp | 11 +-
llvm/lib/IR/Instructions.cpp | 1 +
llvm/lib/IR/Operator.cpp | 2 +
llvm/lib/Transforms/Utils/LowerAtomic.cpp | 21 +-
...mw-elementwise.ll => invalid-atomicrmw.ll} | 8 +
.../Inputs/invalid-atomicrmw-fmf-surplus.bc | Bin 0 -> 2028 bytes
.../Inputs/invalid-atomicrmw-fmf-truncated.bc | Bin 0 -> 2028 bytes
llvm/test/Bitcode/atomic.ll | 19 ++
llvm/test/Bitcode/invalid.test | 7 +
.../GlobalISel/irtranslator-atomicrmw.ll | 4 +-
.../AtomicExpand/AArch64/atomicrmw-fp.ll | 4 +-
.../Transforms/LowerAtomic/atomic-load.ll | 12 +-
.../SelectionDAGNodeConstructionTest.cpp | 27 +++
llvm/unittests/IR/IRBuilderTest.cpp | 20 ++
37 files changed, 396 insertions(+), 219 deletions(-)
rename llvm/test/Assembler/{invalid-atomicrmw-elementwise.ll => invalid-atomicrmw.ll} (81%)
create mode 100644 llvm/test/Bitcode/Inputs/invalid-atomicrmw-fmf-surplus.bc
create mode 100644 llvm/test/Bitcode/Inputs/invalid-atomicrmw-fmf-truncated.bc
diff --git a/clang/test/CodeGen/builtins-nvptx-ptx50.cu b/clang/test/CodeGen/builtins-nvptx-ptx50.cu
index 2a141baf3a6d0..25d1a8575e0a4 100644
--- a/clang/test/CodeGen/builtins-nvptx-ptx50.cu
+++ b/clang/test/CodeGen/builtins-nvptx-ptx50.cu
@@ -17,7 +17,7 @@
// CHECK-LABEL: test_fn
__device__ void test_fn(double d, double* double_ptr) {
- // CHECK: atomicrmw fadd ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw contract fadd ptr {{.*}} monotonic, align 8
// expected-error at +1 {{'__nvvm_atom_add_gen_d' needs target feature sm_60}}
__nvvm_atom_add_gen_d(double_ptr, d);
}
diff --git a/clang/test/CodeGen/fp-atomic-ops.c b/clang/test/CodeGen/fp-atomic-ops.c
index c894e7b4ade37..0e7a2588d42d2 100644
--- a/clang/test/CodeGen/fp-atomic-ops.c
+++ b/clang/test/CodeGen/fp-atomic-ops.c
@@ -19,6 +19,9 @@
// RUN: %clang_cc1 %s -emit-llvm -DDOUBLE -O0 -o - -triple=x86_64-linux-gnu \
// RUN: | FileCheck -check-prefixes=FLOAT,DOUBLE %s
+// RUN: %clang_cc1 %s -emit-llvm -O0 -ffast-math -ffp-contract=fast -o - \
+// RUN: -triple=x86_64-linux-gnu | FileCheck -check-prefix=FAST %s
+
typedef enum memory_order {
memory_order_relaxed = __ATOMIC_RELAXED,
memory_order_acquire = __ATOMIC_ACQUIRE,
@@ -29,9 +32,11 @@ typedef enum memory_order {
void test(float *f, float ff, double *d, double dd) {
// FLOAT: atomicrmw fadd ptr {{.*}} monotonic
+ // FAST: atomicrmw fast fadd ptr {{.*}} monotonic
__atomic_fetch_add(f, ff, memory_order_relaxed);
// FLOAT: atomicrmw fsub ptr {{.*}} monotonic
+ // FAST: atomicrmw fast fsub ptr {{.*}} monotonic
__atomic_fetch_sub(f, ff, memory_order_relaxed);
#ifdef DOUBLE
diff --git a/clang/test/CodeGenCUDA/atomic-options.hip b/clang/test/CodeGenCUDA/atomic-options.hip
index 7b319516a1010..e9afaee3a224b 100644
--- a/clang/test/CodeGenCUDA/atomic-options.hip
+++ b/clang/test/CodeGenCUDA/atomic-options.hip
@@ -25,7 +25,7 @@
// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// HOST-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: ret void
@@ -41,7 +41,7 @@
// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]]
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3:![0-9]+]], !amdgpu.no.remote.memory [[META3]]
// DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: ret void
@@ -57,7 +57,7 @@
// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3:![0-9]+]], !amdgpu.ignore.denormal.mode [[META3]]
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3:![0-9]+]], !amdgpu.ignore.denormal.mode [[META3]]
// OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: ret void
@@ -73,7 +73,7 @@
// SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4:![0-9]+]], !amdgpu.no.remote.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: ret void
@@ -89,7 +89,7 @@
// SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4:![0-9]+]], !amdgpu.ignore.denormal.mode [[META4]]
+// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4:![0-9]+]], !amdgpu.ignore.denormal.mode [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: ret void
@@ -108,7 +108,7 @@ __device__ __host__ void test_default(float *a) {
// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// HOST-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: ret void
@@ -124,7 +124,7 @@ __device__ __host__ void test_default(float *a) {
// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]]
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]]
// DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: ret void
@@ -140,7 +140,7 @@ __device__ __host__ void test_default(float *a) {
// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]]
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]]
// OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: ret void
@@ -156,7 +156,7 @@ __device__ __host__ void test_default(float *a) {
// SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: ret void
@@ -172,7 +172,7 @@ __device__ __host__ void test_default(float *a) {
// SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
+// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: ret void
@@ -193,7 +193,7 @@ __device__ __host__ void test_one(float *a) {
// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// HOST-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: ret void
@@ -209,7 +209,7 @@ __device__ __host__ void test_one(float *a) {
// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]]
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]]
// DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: ret void
@@ -225,7 +225,7 @@ __device__ __host__ void test_one(float *a) {
// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META3]]
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META3]]
// OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: ret void
@@ -241,7 +241,7 @@ __device__ __host__ void test_one(float *a) {
// SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
+// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: ret void
@@ -257,7 +257,7 @@ __device__ __host__ void test_one(float *a) {
// SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]
+// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: ret void
@@ -278,7 +278,7 @@ __device__ __host__ void test_two(float *a) {
// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// HOST-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: ret void
@@ -294,7 +294,7 @@ __device__ __host__ void test_two(float *a) {
// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]]
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]]
// DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: ret void
@@ -310,7 +310,7 @@ __device__ __host__ void test_two(float *a) {
// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]]
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]]
// OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: ret void
@@ -326,7 +326,7 @@ __device__ __host__ void test_two(float *a) {
// SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: ret void
@@ -342,7 +342,7 @@ __device__ __host__ void test_two(float *a) {
// SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]
+// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: ret void
@@ -363,7 +363,7 @@ __device__ __host__ void test_three(float *a) {
// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// HOST-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: ret void
@@ -379,7 +379,7 @@ __device__ __host__ void test_three(float *a) {
// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]]
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]]
// DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: ret void
@@ -395,7 +395,7 @@ __device__ __host__ void test_three(float *a) {
// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META3]]
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META3]]
// OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: ret void
@@ -411,7 +411,7 @@ __device__ __host__ void test_three(float *a) {
// SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: ret void
@@ -427,7 +427,7 @@ __device__ __host__ void test_three(float *a) {
// SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]
+// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.ignore.denormal.mode [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: ret void
@@ -454,25 +454,25 @@ __device__ __host__ void test_multiple_attrs(float *a) {
// HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// HOST-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// HOST-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// HOST-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
// HOST-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// HOST-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 2.000000e+00, ptr [[DOTATOMICTMP1]], align 4
// HOST-NEXT: [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1]], align 4
-// HOST-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] seq_cst, align 4
+// HOST-NEXT: [[TMP6:%.*]] = atomicrmw contract fmax ptr [[TMP4]], float [[TMP5]] seq_cst, align 4
// HOST-NEXT: store float [[TMP6]], ptr [[ATOMIC_TEMP2]], align 4
// HOST-NEXT: [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2]], align 4
// HOST-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 3.000000e+00, ptr [[DOTATOMICTMP3]], align 4
// HOST-NEXT: [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3]], align 4
-// HOST-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] acquire, align 4
+// HOST-NEXT: [[TMP10:%.*]] = atomicrmw contract fmin ptr [[TMP8]], float [[TMP9]] acquire, align 4
// HOST-NEXT: store float [[TMP10]], ptr [[ATOMIC_TEMP4]], align 4
// HOST-NEXT: [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4]], align 4
// HOST-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR]], align 8
// HOST-NEXT: store float 4.000000e+00, ptr [[DOTATOMICTMP5]], align 4
// HOST-NEXT: [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5]], align 4
-// HOST-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] release, align 4
+// HOST-NEXT: [[TMP14:%.*]] = atomicrmw contract fsub ptr [[TMP12]], float [[TMP13]] release, align 4
// HOST-NEXT: store float [[TMP14]], ptr [[ATOMIC_TEMP6]], align 4
// HOST-NEXT: [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6]], align 4
// HOST-NEXT: ret void
@@ -494,25 +494,25 @@ __device__ __host__ void test_multiple_attrs(float *a) {
// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// DEV-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]]
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META3]], !amdgpu.no.remote.memory [[META3]]
// DEV-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// DEV-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 2.000000e+00, ptr addrspace(5) [[DOTATOMICTMP1]], align 4
// DEV-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP1]], align 4
-// DEV-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4
+// DEV-NEXT: [[TMP6:%.*]] = atomicrmw contract fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4
// DEV-NEXT: store float [[TMP6]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4
// DEV-NEXT: [[TMP7:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4
// DEV-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 3.000000e+00, ptr addrspace(5) [[DOTATOMICTMP3]], align 4
// DEV-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP3]], align 4
-// DEV-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META3]]
+// DEV-NEXT: [[TMP10:%.*]] = atomicrmw contract fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META3]]
// DEV-NEXT: store float [[TMP10]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4
// DEV-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4
// DEV-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// DEV-NEXT: store float 4.000000e+00, ptr addrspace(5) [[DOTATOMICTMP5]], align 4
// DEV-NEXT: [[TMP13:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP5]], align 4
-// DEV-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META3]]
+// DEV-NEXT: [[TMP14:%.*]] = atomicrmw contract fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META3]]
// DEV-NEXT: store float [[TMP14]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4
// DEV-NEXT: [[TMP15:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4
// DEV-NEXT: ret void
@@ -534,25 +534,25 @@ __device__ __host__ void test_multiple_attrs(float *a) {
// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 1.000000e+00, ptr addrspace(5) [[DOTATOMICTMP]], align 4
// OPT-NEXT: [[TMP1:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP]], align 4
-// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]]
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META3]], !amdgpu.ignore.denormal.mode [[META3]]
// OPT-NEXT: store float [[TMP2]], ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP3:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP]], align 4
// OPT-NEXT: [[TMP4:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 2.000000e+00, ptr addrspace(5) [[DOTATOMICTMP1]], align 4
// OPT-NEXT: [[TMP5:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP1]], align 4
-// OPT-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4
+// OPT-NEXT: [[TMP6:%.*]] = atomicrmw contract fmax ptr [[TMP4]], float [[TMP5]] syncscope("agent") seq_cst, align 4
// OPT-NEXT: store float [[TMP6]], ptr addrspace(5) [[ATOMIC_TEMP2]], align 4
// OPT-NEXT: [[TMP7:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP2]], align 4
// OPT-NEXT: [[TMP8:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 3.000000e+00, ptr addrspace(5) [[DOTATOMICTMP3]], align 4
// OPT-NEXT: [[TMP9:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP3]], align 4
-// OPT-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META3]]
+// OPT-NEXT: [[TMP10:%.*]] = atomicrmw contract fmin ptr [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META3]]
// OPT-NEXT: store float [[TMP10]], ptr addrspace(5) [[ATOMIC_TEMP4]], align 4
// OPT-NEXT: [[TMP11:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP4]], align 4
// OPT-NEXT: [[TMP12:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
// OPT-NEXT: store float 4.000000e+00, ptr addrspace(5) [[DOTATOMICTMP5]], align 4
// OPT-NEXT: [[TMP13:%.*]] = load float, ptr addrspace(5) [[DOTATOMICTMP5]], align 4
-// OPT-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META3]]
+// OPT-NEXT: [[TMP14:%.*]] = atomicrmw contract fsub ptr [[TMP12]], float [[TMP13]] syncscope("wavefront") release, align 4, !amdgpu.no.fine.grained.memory [[META3]]
// OPT-NEXT: store float [[TMP14]], ptr addrspace(5) [[ATOMIC_TEMP6]], align 4
// OPT-NEXT: [[TMP15:%.*]] = load float, ptr addrspace(5) [[ATOMIC_TEMP6]], align 4
// OPT-NEXT: ret void
@@ -574,25 +574,25 @@ __device__ __host__ void test_multiple_attrs(float *a) {
// SPIRV-DEV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.fine.grained.memory [[META4]], !amdgpu.no.remote.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-DEV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 2.000000e+00, ptr [[DOTATOMICTMP1]], align 4
// SPIRV-DEV-NEXT: [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1]], align 4
-// SPIRV-DEV-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr addrspace(4) [[TMP4]], float [[TMP5]] syncscope("device") seq_cst, align 4
+// SPIRV-DEV-NEXT: [[TMP6:%.*]] = atomicrmw contract fmax ptr addrspace(4) [[TMP4]], float [[TMP5]] syncscope("device") seq_cst, align 4
// SPIRV-DEV-NEXT: store float [[TMP6]], ptr [[ATOMIC_TEMP2]], align 4
// SPIRV-DEV-NEXT: [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2]], align 4
// SPIRV-DEV-NEXT: [[TMP8:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 3.000000e+00, ptr [[DOTATOMICTMP3]], align 4
// SPIRV-DEV-NEXT: [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3]], align 4
-// SPIRV-DEV-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr addrspace(4) [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP10:%.*]] = atomicrmw contract fmin ptr addrspace(4) [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP10]], ptr [[ATOMIC_TEMP4]], align 4
// SPIRV-DEV-NEXT: [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4]], align 4
// SPIRV-DEV-NEXT: [[TMP12:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-DEV-NEXT: store float 4.000000e+00, ptr [[DOTATOMICTMP5]], align 4
// SPIRV-DEV-NEXT: [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5]], align 4
-// SPIRV-DEV-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr addrspace(4) [[TMP12]], float [[TMP13]] syncscope("subgroup") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]
+// SPIRV-DEV-NEXT: [[TMP14:%.*]] = atomicrmw contract fsub ptr addrspace(4) [[TMP12]], float [[TMP13]] syncscope("subgroup") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]
// SPIRV-DEV-NEXT: store float [[TMP14]], ptr [[ATOMIC_TEMP6]], align 4
// SPIRV-DEV-NEXT: [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6]], align 4
// SPIRV-DEV-NEXT: ret void
@@ -614,25 +614,25 @@ __device__ __host__ void test_multiple_attrs(float *a) {
// SPIRV-OPT-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
// SPIRV-OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
-// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
+// SPIRV-OPT-NEXT: [[TMP2:%.*]] = atomicrmw contract fadd ptr addrspace(4) [[TMP0]], float [[TMP1]] monotonic, align 4, !amdgpu.no.remote.memory [[META4]], !amdgpu.ignore.denormal.mode [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
// SPIRV-OPT-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 2.000000e+00, ptr [[DOTATOMICTMP1]], align 4
// SPIRV-OPT-NEXT: [[TMP5:%.*]] = load float, ptr [[DOTATOMICTMP1]], align 4
-// SPIRV-OPT-NEXT: [[TMP6:%.*]] = atomicrmw fmax ptr addrspace(4) [[TMP4]], float [[TMP5]] syncscope("device") seq_cst, align 4
+// SPIRV-OPT-NEXT: [[TMP6:%.*]] = atomicrmw contract fmax ptr addrspace(4) [[TMP4]], float [[TMP5]] syncscope("device") seq_cst, align 4
// SPIRV-OPT-NEXT: store float [[TMP6]], ptr [[ATOMIC_TEMP2]], align 4
// SPIRV-OPT-NEXT: [[TMP7:%.*]] = load float, ptr [[ATOMIC_TEMP2]], align 4
// SPIRV-OPT-NEXT: [[TMP8:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 3.000000e+00, ptr [[DOTATOMICTMP3]], align 4
// SPIRV-OPT-NEXT: [[TMP9:%.*]] = load float, ptr [[DOTATOMICTMP3]], align 4
-// SPIRV-OPT-NEXT: [[TMP10:%.*]] = atomicrmw fmin ptr addrspace(4) [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]
+// SPIRV-OPT-NEXT: [[TMP10:%.*]] = atomicrmw contract fmin ptr addrspace(4) [[TMP8]], float [[TMP9]] syncscope("workgroup") acquire, align 4, !amdgpu.no.remote.memory [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP10]], ptr [[ATOMIC_TEMP4]], align 4
// SPIRV-OPT-NEXT: [[TMP11:%.*]] = load float, ptr [[ATOMIC_TEMP4]], align 4
// SPIRV-OPT-NEXT: [[TMP12:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
// SPIRV-OPT-NEXT: store float 4.000000e+00, ptr [[DOTATOMICTMP5]], align 4
// SPIRV-OPT-NEXT: [[TMP13:%.*]] = load float, ptr [[DOTATOMICTMP5]], align 4
-// SPIRV-OPT-NEXT: [[TMP14:%.*]] = atomicrmw fsub ptr addrspace(4) [[TMP12]], float [[TMP13]] syncscope("subgroup") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]
+// SPIRV-OPT-NEXT: [[TMP14:%.*]] = atomicrmw contract fsub ptr addrspace(4) [[TMP12]], float [[TMP13]] syncscope("subgroup") release, align 4, !amdgpu.no.fine.grained.memory [[META4]]
// SPIRV-OPT-NEXT: store float [[TMP14]], ptr [[ATOMIC_TEMP6]], align 4
// SPIRV-OPT-NEXT: [[TMP15:%.*]] = load float, ptr [[ATOMIC_TEMP6]], align 4
// SPIRV-OPT-NEXT: ret void
diff --git a/clang/test/CodeGenCUDA/builtins-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-amdgcn.cu
index 47c6ba57ec2f2..4bf4a3843a6d5 100644
--- a/clang/test/CodeGenCUDA/builtins-amdgcn.cu
+++ b/clang/test/CodeGenCUDA/builtins-amdgcn.cu
@@ -92,7 +92,7 @@ __global__
// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
// CHECK-NEXT: store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 4
+// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw contract fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP1]], ptr [[X_ASCAST]], align 4
// CHECK-NEXT: ret void
//
@@ -109,7 +109,7 @@ __global__
// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
// CHECK-NEXT: store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4
+// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw contract fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP1]], ptr [[X_ASCAST]], align 4
// CHECK-NEXT: ret void
//
@@ -134,7 +134,7 @@ __global__ void test_ds_fadd(float src) {
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw contract fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP3]], ptr [[X_ASCAST]], align 4
// CHECK-NEXT: ret void
//
@@ -232,7 +232,7 @@ __device__ void func(float *x);
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw contract fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP3]], ptr [[X_ASCAST]], align 4
// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8
// CHECK-NEXT: call void @_Z4funcPf(ptr noundef [[TMP4]]) #[[ATTR8:[0-9]+]]
diff --git a/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu
index ccd76037b03cf..e3e1bffa910ed 100644
--- a/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu
+++ b/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu
@@ -140,7 +140,7 @@ __global__ void use_implicitarg_ptr(int* out) {
// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)
// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 4
+// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw contract fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP1]], ptr addrspace(4) [[X_ASCAST]], align 4
// CHECK-NEXT: ret void
//
@@ -152,7 +152,7 @@ __global__ void use_implicitarg_ptr(int* out) {
// AMDGCNSPIRV-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)
// AMDGCNSPIRV-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = atomicrmw fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 4
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = atomicrmw contract fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 4
// AMDGCNSPIRV-NEXT: store volatile float [[TMP1]], ptr addrspace(4) [[X_ASCAST]], align 4
// AMDGCNSPIRV-NEXT: ret void
//
@@ -171,7 +171,7 @@ __global__ void test_ds_fmax(float src) {
// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)
// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4
+// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw contract fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP1]], ptr addrspace(4) [[X_ASCAST]], align 4
// CHECK-NEXT: ret void
//
@@ -183,7 +183,7 @@ __global__ void test_ds_fmax(float src) {
// AMDGCNSPIRV-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)
// AMDGCNSPIRV-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = atomicrmw fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4
+// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = atomicrmw contract fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4
// AMDGCNSPIRV-NEXT: store volatile float [[TMP1]], ptr addrspace(4) [[X_ASCAST]], align 4
// AMDGCNSPIRV-NEXT: ret void
//
@@ -208,7 +208,7 @@ __global__ void test_ds_fadd(float src) {
// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw contract fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP3]], ptr addrspace(4) [[X_ASCAST]], align 4
// CHECK-NEXT: ret void
//
@@ -228,7 +228,7 @@ __global__ void test_ds_fadd(float src) {
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)
// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = atomicrmw contract fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// AMDGCNSPIRV-NEXT: store volatile float [[TMP3]], ptr addrspace(4) [[X_ASCAST]], align 4
// AMDGCNSPIRV-NEXT: ret void
//
@@ -371,7 +371,7 @@ __device__ void func(float *x);
// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw contract fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// CHECK-NEXT: store volatile float [[TMP3]], ptr addrspace(4) [[X_ASCAST]], align 4
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8
// CHECK-NEXT: call spir_func addrspace(4) void @_Z4funcPf(ptr addrspace(4) noundef [[TMP4]]) #[[ATTR8:[0-9]+]]
@@ -393,7 +393,7 @@ __device__ void func(float *x);
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8
// AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)
// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4
-// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = atomicrmw contract fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// AMDGCNSPIRV-NEXT: store volatile float [[TMP3]], ptr addrspace(4) [[X_ASCAST]], align 4
// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8
// AMDGCNSPIRV-NEXT: call spir_func addrspace(4) void @_Z4funcPf(ptr addrspace(4) noundef [[TMP4]]) #[[ATTR8:[0-9]+]]
diff --git a/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu b/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
index 03b39cd1291c7..1320a4a38e551 100644
--- a/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
+++ b/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
@@ -12,7 +12,7 @@ typedef __attribute__((address_space(3))) float *LP;
// CHECK: %[[ADDR_ADDR_ASCAST:.*]] = load ptr, ptr %[[ADDR_ADDR_ASCAST_PTR]], align 8
// CHECK: %[[AS_CAST:.*]] = addrspacecast ptr %[[ADDR_ADDR_ASCAST]] to ptr addrspace(3)
// CHECK: [[TMP2:%.+]] = load float, ptr %val.addr.ascast, align 4
-// CHECK: [[TMP3:%.+]] = atomicrmw fadd ptr addrspace(3) %[[AS_CAST]], float [[TMP2]] monotonic, align 4
+// CHECK: [[TMP3:%.+]] = atomicrmw contract fadd ptr addrspace(3) %[[AS_CAST]], float [[TMP2]] monotonic, align 4
// CHECK: %4 = load ptr, ptr %rtn.ascast, align 8
// CHECK: store float [[TMP3]], ptr %4, align 4
__device__ void test_ds_atomic_add_f32(float *addr, float val) {
diff --git a/clang/test/CodeGenCUDA/builtins-unsafe-atomics-spirv-amdgcn-gfx90a.cu b/clang/test/CodeGenCUDA/builtins-unsafe-atomics-spirv-amdgcn-gfx90a.cu
index e01d9d7efed27..315ab19ad4238 100644
--- a/clang/test/CodeGenCUDA/builtins-unsafe-atomics-spirv-amdgcn-gfx90a.cu
+++ b/clang/test/CodeGenCUDA/builtins-unsafe-atomics-spirv-amdgcn-gfx90a.cu
@@ -20,7 +20,7 @@ typedef __attribute__((address_space(3))) float *LP;
// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ADDR_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 4
-// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fadd ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
+// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw contract fadd ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[RTN_ASCAST]], align 8
// CHECK-NEXT: store float [[TMP3]], ptr addrspace(4) [[TMP4]], align 4
// CHECK-NEXT: ret void
diff --git a/clang/test/SemaHIP/incorrect-atomic-scope.hip b/clang/test/SemaHIP/incorrect-atomic-scope.hip
index 1c5aaee710051..5aa48162e8875 100644
--- a/clang/test/SemaHIP/incorrect-atomic-scope.hip
+++ b/clang/test/SemaHIP/incorrect-atomic-scope.hip
@@ -7,7 +7,7 @@
// is generated.
//
// CHECK-LABEL: test_builtin_rmw
-// CHECK: atomicrmw fmax {{.*}} syncscope("singlethread")
+// CHECK: atomicrmw contract fmax {{.*}} syncscope("singlethread")
//
// CHECK-LABEL: test_scoped_atomic
// CHECK: atomicrmw {{.*}} syncscope("singlethread")
diff --git a/llvm/docs/LangRef.md b/llvm/docs/LangRef.md
index 5275cf87d7cad..db0a2218ecd74 100644
--- a/llvm/docs/LangRef.md
+++ b/llvm/docs/LangRef.md
@@ -4307,9 +4307,10 @@ LLVM IR floating-point operations ({ref}`fneg <i_fneg>`, {ref}`fadd <i_fadd>`,
{ref}`fsub <i_fsub>`, {ref}`fmul <i_fmul>`, {ref}`fdiv <i_fdiv>`,
{ref}`frem <i_frem>`, {ref}`fcmp <i_fcmp>`, {ref}`fptrunc <i_fptrunc>`,
{ref}`fpext <i_fpext>`), {ref}`uitofp <i_uitofp>`, {ref}`sitofp <i_sitofp>`,
-and {ref}`phi <i_phi>`, {ref}`select <i_select>`, or {ref}`call <i_call>`
-instructions that return floating-point types may use the following flags to
-enable otherwise unsafe floating-point transformations.
+and {ref}`phi <i_phi>`, {ref}`select <i_select>`, {ref}`call <i_call>` or
+{ref}`atomicrmw <i_atomicrmw>` instructions that return floating-point types,
+may use the following flags to enable otherwise unsafe floating-point
+transformations.
`fast`
: This flag is a shorthand for specifying all fast-math flags at once, and
@@ -4340,6 +4341,14 @@ types:
- Array types (nested to any depth) of floating-point scalar or vector types
- Homogeneous literal struct types of floating-point scalar or vector types
+(fastmath_atomicrmw_types)=
+
+For {ref}`atomicrmw <i_atomicrmw>` instructions, the following types are
+considered to be floating-point types:
+
+- Floating-point scalar types
+- Fixed vector types of floating-point elements
+
#### Rewrite-based flags
The following flags have rewrite-based semantics. These flags allow expressions,
@@ -12085,7 +12094,7 @@ done:
##### Syntax:
```
-atomicrmw [volatile] [elementwise] <operation> ptr <pointer>, <ty> <value> [syncscope("<target-scope>")] <ordering>[, align <alignment>] ; yields ty
+atomicrmw [volatile] [fast-math flags]* [elementwise] <operation> ptr <pointer>, <ty> <value> [syncscope("<target-scope>")] <ordering>[, align <alignment>] ; yields ty
```
##### Overview:
@@ -12127,6 +12136,11 @@ For add/sub/and/nand/or/xor/max/min/umax/umin/uinc_wrap/udec_wrap/usub_cond/usub
For fadd/fsub/fmax/fmin/fmaximum/fminimum/fmaximumnum/fminimumnum, this must be a floating-point or fixed vector of floating-point type.
For xchg, this must be an integer type, floating-point type, or pointer type, or, if the `elementwise` modifier is present, a fixed vector of integer type, floating-point type, or pointer type.
The type of the `<pointer>` operand must be a pointer to the type of `<value>`.
+An `atomicrmw` can also take any number of {ref}`fast-math flags <fastmath>`,
+which are optimization hints to enable otherwise unsafe floating-point
+optimizations. Fast-math flags are only valid for `atomicrmw` instructions
+operating on {ref}`supported floating-point types <fastmath_atomicrmw_types>`.
+The flags additionally apply to the value written to memory.
If the `atomicrmw` is marked as `volatile`, then the optimizer is not allowed to modify the
number or order of execution of this `atomicrmw` with other
{ref}`volatile operations <volatile>`.
diff --git a/llvm/include/llvm/Bitcode/LLVMBitCodes.h b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
index 358f9a65a80af..d875b5e8dd9d3 100644
--- a/llvm/include/llvm/Bitcode/LLVMBitCodes.h
+++ b/llvm/include/llvm/Bitcode/LLVMBitCodes.h
@@ -531,6 +531,7 @@ enum RMWOperations {
enum RMWOperationFlags {
RMW_ELEMENTWISE_FLAG = 1 << 5,
+ RMW_FMF_FLAG = 1 << 6,
};
/// OverflowingBinaryOperatorOptionalFlags - Flags for serializing
@@ -692,8 +693,8 @@ enum FunctionCodes {
// fnty, fnid, args...]
FUNC_CODE_INST_FREEZE = 58, // FREEZE: [opty, opval]
FUNC_CODE_INST_ATOMICRMW = 59, // ATOMICRMW: [ptrty, ptr, valty, val,
- // operation, align, vol,
- // ordering, syncscope]
+ // operation, fmf?, vol, ordering,
+ // syncscope, align]
FUNC_CODE_BLOCKADDR_USERS = 60, // BLOCKADDR_USERS: [value...]
FUNC_CODE_DEBUG_RECORD_VALUE =
diff --git a/llvm/include/llvm/CodeGen/GlobalISel/MachineIRBuilder.h b/llvm/include/llvm/CodeGen/GlobalISel/MachineIRBuilder.h
index ee533d2a8d8b8..be83701d7f350 100644
--- a/llvm/include/llvm/CodeGen/GlobalISel/MachineIRBuilder.h
+++ b/llvm/include/llvm/CodeGen/GlobalISel/MachineIRBuilder.h
@@ -1543,7 +1543,9 @@ class LLVM_ABI MachineIRBuilder {
/// \return a MachineInstrBuilder for the newly created instruction.
MachineInstrBuilder buildAtomicRMW(unsigned Opcode, const DstOp &OldValRes,
const SrcOp &Addr, const SrcOp &Val,
- MachineMemOperand &MMO);
+ MachineMemOperand &MMO,
+ std::optional<unsigned> Flags =
+ std::nullopt);
/// Build and insert `OldValRes<def> = G_ATOMICRMW_XCHG Addr, Val, MMO`.
///
diff --git a/llvm/include/llvm/CodeGen/SelectionDAG.h b/llvm/include/llvm/CodeGen/SelectionDAG.h
index abc2e1f0ebe4f..c1095f63de49f 100644
--- a/llvm/include/llvm/CodeGen/SelectionDAG.h
+++ b/llvm/include/llvm/CodeGen/SelectionDAG.h
@@ -1453,14 +1453,16 @@ class SelectionDAG {
/// and chain and takes 2 operands.
LLVM_ABI SDValue getAtomic(unsigned Opcode, const SDLoc &dl, EVT MemVT,
SDValue Chain, SDValue Ptr, SDValue Val,
- MachineMemOperand *MMO);
+ MachineMemOperand *MMO,
+ const SDNodeFlags Flags = {});
/// Gets a node for an atomic op, produces result and chain and takes N
/// operands.
LLVM_ABI SDValue getAtomic(unsigned Opcode, const SDLoc &dl, EVT MemVT,
SDVTList VTList, ArrayRef<SDValue> Ops,
MachineMemOperand *MMO,
- ISD::LoadExtType ExtType = ISD::NON_EXTLOAD);
+ ISD::LoadExtType ExtType = ISD::NON_EXTLOAD,
+ const SDNodeFlags Flags = {});
LLVM_ABI SDValue getAtomicLoad(ISD::LoadExtType ExtType, const SDLoc &dl,
EVT MemVT, EVT VT, SDValue Chain, SDValue Ptr,
diff --git a/llvm/include/llvm/IR/IRBuilder.h b/llvm/include/llvm/IR/IRBuilder.h
index f2621bbe298df..abf181ab2dee7 100644
--- a/llvm/include/llvm/IR/IRBuilder.h
+++ b/llvm/include/llvm/IR/IRBuilder.h
@@ -1052,24 +1052,28 @@ class IRBuilderBase {
}
/// Create call to the minimum intrinsic.
- Value *CreateMinimum(Value *LHS, Value *RHS, const Twine &Name = "") {
- return CreateBinaryIntrinsic(Intrinsic::minimum, LHS, RHS, nullptr, Name);
+ Value *CreateMinimum(Value *LHS, Value *RHS, FMFSource FMFSource = {},
+ const Twine &Name = "") {
+ return CreateBinaryIntrinsic(Intrinsic::minimum, LHS, RHS, FMFSource, Name);
}
/// Create call to the maximum intrinsic.
- Value *CreateMaximum(Value *LHS, Value *RHS, const Twine &Name = "") {
- return CreateBinaryIntrinsic(Intrinsic::maximum, LHS, RHS, nullptr, Name);
+ Value *CreateMaximum(Value *LHS, Value *RHS, FMFSource FMFSource = {},
+ const Twine &Name = "") {
+ return CreateBinaryIntrinsic(Intrinsic::maximum, LHS, RHS, FMFSource, Name);
}
/// Create call to the minimumnum intrinsic.
- Value *CreateMinimumNum(Value *LHS, Value *RHS, const Twine &Name = "") {
- return CreateBinaryIntrinsic(Intrinsic::minimumnum, LHS, RHS, nullptr,
+ Value *CreateMinimumNum(Value *LHS, Value *RHS, FMFSource FMFSource = {},
+ const Twine &Name = "") {
+ return CreateBinaryIntrinsic(Intrinsic::minimumnum, LHS, RHS, FMFSource,
Name);
}
/// Create call to the maximum intrinsic.
- Value *CreateMaximumNum(Value *LHS, Value *RHS, const Twine &Name = "") {
- return CreateBinaryIntrinsic(Intrinsic::maximumnum, LHS, RHS, nullptr,
+ Value *CreateMaximumNum(Value *LHS, Value *RHS, FMFSource FMFSource = {},
+ const Twine &Name = "") {
+ return CreateBinaryIntrinsic(Intrinsic::maximumnum, LHS, RHS, FMFSource,
Name);
}
@@ -1988,8 +1992,11 @@ class IRBuilderBase {
Align = llvm::Align(DL.getTypeStoreSize(Val->getType()));
}
- return Insert(
- new AtomicRMWInst(Op, Ptr, Val, *Align, Ordering, SSID, Elementwise));
+ AtomicRMWInst *RMWI =
+ new AtomicRMWInst(Op, Ptr, Val, *Align, Ordering, SSID, Elementwise);
+ if (isa<FPMathOperator>(RMWI))
+ setFPAttrs(RMWI, nullptr, FMF);
+ return Insert(RMWI);
}
Value *CreateStructuredGEP(Type *BaseType, Value *PtrBase,
diff --git a/llvm/include/llvm/IR/Instructions.h b/llvm/include/llvm/IR/Instructions.h
index 5838109847845..679f84ff8bc25 100644
--- a/llvm/include/llvm/IR/Instructions.h
+++ b/llvm/include/llvm/IR/Instructions.h
@@ -757,7 +757,7 @@ DEFINE_TRANSPARENT_OPERAND_ACCESSORS(AtomicCmpXchgInst, Value)
/// combines it with another value, and then stores the result back. Returns
/// the old value.
///
-class AtomicRMWInst : public Instruction {
+class AtomicRMWInst : public Instruction, public FastMathFlagsStorage {
protected:
// Note: Instruction needs to be a friend here to call cloneImpl.
friend class Instruction;
diff --git a/llvm/include/llvm/IR/Operator.h b/llvm/include/llvm/IR/Operator.h
index c44cddca93009..bbd8eb40f1149 100644
--- a/llvm/include/llvm/IR/Operator.h
+++ b/llvm/include/llvm/IR/Operator.h
@@ -331,9 +331,10 @@ class FPMathOperator : public Operator {
return true;
case Instruction::PHI:
case Instruction::Select:
- case Instruction::Call: {
+ case Instruction::Call:
return isSupportedFloatingPointType(V->getType());
- }
+ case Instruction::AtomicRMW:
+ return V->getType()->isFPOrFPVectorTy();
default:
return false;
}
diff --git a/llvm/include/llvm/Transforms/Utils/LowerAtomic.h b/llvm/include/llvm/Transforms/Utils/LowerAtomic.h
index 44c89fc408684..4be9c0bfe135e 100644
--- a/llvm/include/llvm/Transforms/Utils/LowerAtomic.h
+++ b/llvm/include/llvm/Transforms/Utils/LowerAtomic.h
@@ -39,7 +39,7 @@ LLVM_ABI bool lowerAtomicRMWInst(AtomicRMWInst *RMWI);
/// returning the new value.
LLVM_ABI Value *buildAtomicRMWValue(AtomicRMWInst::BinOp Op,
IRBuilderBase &Builder, Value *Loaded,
- Value *Val);
+ Value *Val, FastMathFlags FMF = {});
}
#endif // LLVM_TRANSFORMS_UTILS_LOWERATOMIC_H
diff --git a/llvm/lib/AsmParser/LLParser.cpp b/llvm/lib/AsmParser/LLParser.cpp
index c58b109b5ff9b..dea7382f576dd 100644
--- a/llvm/lib/AsmParser/LLParser.cpp
+++ b/llvm/lib/AsmParser/LLParser.cpp
@@ -9133,8 +9133,8 @@ int LLParser::parseCmpXchg(Instruction *&Inst, PerFunctionState &PFS) {
}
/// parseAtomicRMW
-/// ::= 'atomicrmw' 'volatile'? 'elementwise'? BinOp TypeAndValue ','
-/// TypeAndValue
+/// ::= 'atomicrmw' 'volatile'? 'fast-math flags'? 'elementwise'? BinOp
+/// TypeAndValue ',' TypeAndValue
/// 'singlethread'? AtomicOrdering
int LLParser::parseAtomicRMW(Instruction *&Inst, PerFunctionState &PFS) {
Value *Ptr, *Val; LocTy PtrLoc, ValLoc;
@@ -9149,6 +9149,7 @@ int LLParser::parseAtomicRMW(Instruction *&Inst, PerFunctionState &PFS) {
if (EatIfPresent(lltok::kw_volatile))
IsVolatile = true;
+ FastMathFlags FMF = EatFastMathFlagsIfPresent();
if (EatIfPresent(lltok::kw_elementwise))
IsElementwise = true;
@@ -9268,12 +9269,17 @@ int LLParser::parseAtomicRMW(Instruction *&Inst, PerFunctionState &PFS) {
if (Size < 8 || (Size & (Size - 1)))
return error(ValLoc,
"atomicrmw operand must have a power-of-two byte size");
+ if (FMF.any() && !Val->getType()->isFPOrFPVectorTy())
+ return error(ValLoc, "fast-math-flags specified for atomicrmw without "
+ "floating-point type");
const Align DefaultAlignment(
PFS.getFunction().getDataLayout().getTypeStoreSize(Val->getType()));
AtomicRMWInst *RMWI = new AtomicRMWInst(Operation, Ptr, Val,
Alignment.value_or(DefaultAlignment),
Ordering, SSID, IsElementwise);
RMWI->setVolatile(IsVolatile);
+ if (FMF.any())
+ RMWI->setFastMathFlags(FMF);
Inst = RMWI;
return AteExtraComma ? InstExtraComma : InstNormal;
}
diff --git a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
index ac61ede6395af..befdc5dd332ef 100644
--- a/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
+++ b/llvm/lib/Bitcode/Reader/BitcodeReader.cpp
@@ -1395,10 +1395,12 @@ static int getDecodedBinaryOpcode(unsigned Val, Type *Ty) {
}
}
-static AtomicRMWInst::BinOp getDecodedRMWOperation(unsigned Val,
- bool &IsElementwise) {
+static AtomicRMWInst::BinOp
+getDecodedRMWOperation(unsigned Val, bool &IsElementwise,
+ bool &HasFastMathFlags) {
IsElementwise = Val & bitc::RMW_ELEMENTWISE_FLAG;
- switch (Val & ~bitc::RMW_ELEMENTWISE_FLAG) {
+ HasFastMathFlags = Val & bitc::RMW_FMF_FLAG;
+ switch (Val & ~(bitc::RMW_ELEMENTWISE_FLAG | bitc::RMW_FMF_FLAG)) {
default: return AtomicRMWInst::BAD_BINOP;
case bitc::RMW_XCHG: return AtomicRMWInst::Xchg;
case bitc::RMW_ADD: return AtomicRMWInst::Add;
@@ -6718,7 +6720,8 @@ Error BitcodeReader::parseFunctionBody(Function *F) {
case bitc::FUNC_CODE_INST_ATOMICRMW_OLD:
case bitc::FUNC_CODE_INST_ATOMICRMW: {
// ATOMICRMW_OLD: [ptrty, ptr, val, op, vol, ordering, ssid, align?]
- // ATOMICRMW: [ptrty, ptr, valty, val, op, vol, ordering, ssid, align?]
+ // ATOMICRMW: [ptrty, ptr, valty, val, op, fmf?, vol, ordering, ssid,
+ // align?]
const size_t NumRecords = Record.size();
unsigned OpNum = 0;
@@ -6742,29 +6745,46 @@ Error BitcodeReader::parseFunctionBody(Function *F) {
return error("Invalid atomicrmw record");
}
- if (!(NumRecords == (OpNum + 4) || NumRecords == (OpNum + 5)))
+ if (NumRecords <= OpNum)
return error("Invalid atomicrmw record");
bool IsElementwise = false;
+ bool HasFastMathFlags = false;
const AtomicRMWInst::BinOp Operation =
- getDecodedRMWOperation(Record[OpNum], IsElementwise);
+ getDecodedRMWOperation(Record[OpNum++], IsElementwise,
+ HasFastMathFlags);
if (Operation < AtomicRMWInst::FIRST_BINOP ||
Operation > AtomicRMWInst::LAST_BINOP)
return error("Invalid atomicrmw record");
- const bool IsVol = Record[OpNum + 1];
+ const size_t MinNumRecords = OpNum + (HasFastMathFlags ? 4 : 3);
+ if (!(NumRecords == MinNumRecords || NumRecords == MinNumRecords + 1))
+ return error("Invalid atomicrmw record");
- const AtomicOrdering Ordering = getDecodedOrdering(Record[OpNum + 2]);
+ FastMathFlags FMF;
+ if (HasFastMathFlags) {
+ FMF = getDecodedFastMathFlags(Record[OpNum++]);
+ if (!FMF.any())
+ return error(
+ "Fast math flags indicator set for atomicrmw with no FMF");
+ if (!Val->getType()->isFPOrFPVectorTy())
+ return error("Fast-math-flags specified for atomicrmw without "
+ "floating-point scalar or vector type");
+ }
+
+ const bool IsVol = Record[OpNum++];
+
+ const AtomicOrdering Ordering = getDecodedOrdering(Record[OpNum++]);
if (Ordering == AtomicOrdering::NotAtomic ||
Ordering == AtomicOrdering::Unordered)
return error("Invalid atomicrmw record");
- const SyncScope::ID SSID = getDecodedSyncScopeID(Record[OpNum + 3]);
+ const SyncScope::ID SSID = getDecodedSyncScopeID(Record[OpNum++]);
MaybeAlign Alignment;
- if (NumRecords == (OpNum + 5)) {
- if (Error Err = parseAlignmentValue(Record[OpNum + 4], Alignment))
+ if (NumRecords == (OpNum + 1)) {
+ if (Error Err = parseAlignmentValue(Record[OpNum], Alignment))
return Err;
}
@@ -6776,6 +6796,8 @@ Error BitcodeReader::parseFunctionBody(Function *F) {
IsElementwise);
ResTypeID = ValTypeID;
cast<AtomicRMWInst>(I)->setVolatile(IsVol);
+ if (FMF.any())
+ I->setFastMathFlags(FMF);
InstructionList.push_back(I);
break;
diff --git a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
index 571336c217797..db1680d67d5b2 100644
--- a/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
+++ b/llvm/lib/Bitcode/Writer/BitcodeWriter.cpp
@@ -694,86 +694,6 @@ static unsigned getEncodedBinaryOpcode(unsigned Opcode) {
}
}
-static unsigned getEncodedRMWOperation(const AtomicRMWInst &I) {
- unsigned Encoding = 0;
- switch (I.getOperation()) {
- default: llvm_unreachable("Unknown RMW operation!");
- case AtomicRMWInst::Xchg:
- Encoding = bitc::RMW_XCHG;
- break;
- case AtomicRMWInst::Add:
- Encoding = bitc::RMW_ADD;
- break;
- case AtomicRMWInst::Sub:
- Encoding = bitc::RMW_SUB;
- break;
- case AtomicRMWInst::And:
- Encoding = bitc::RMW_AND;
- break;
- case AtomicRMWInst::Nand:
- Encoding = bitc::RMW_NAND;
- break;
- case AtomicRMWInst::Or:
- Encoding = bitc::RMW_OR;
- break;
- case AtomicRMWInst::Xor:
- Encoding = bitc::RMW_XOR;
- break;
- case AtomicRMWInst::Max:
- Encoding = bitc::RMW_MAX;
- break;
- case AtomicRMWInst::Min:
- Encoding = bitc::RMW_MIN;
- break;
- case AtomicRMWInst::UMax:
- Encoding = bitc::RMW_UMAX;
- break;
- case AtomicRMWInst::UMin:
- Encoding = bitc::RMW_UMIN;
- break;
- case AtomicRMWInst::FAdd:
- Encoding = bitc::RMW_FADD;
- break;
- case AtomicRMWInst::FSub:
- Encoding = bitc::RMW_FSUB;
- break;
- case AtomicRMWInst::FMax:
- Encoding = bitc::RMW_FMAX;
- break;
- case AtomicRMWInst::FMin:
- Encoding = bitc::RMW_FMIN;
- break;
- case AtomicRMWInst::FMaximum:
- Encoding = bitc::RMW_FMAXIMUM;
- break;
- case AtomicRMWInst::FMinimum:
- Encoding = bitc::RMW_FMINIMUM;
- break;
- case AtomicRMWInst::FMaximumNum:
- Encoding = bitc::RMW_FMAXIMUMNUM;
- break;
- case AtomicRMWInst::FMinimumNum:
- Encoding = bitc::RMW_FMINIMUMNUM;
- break;
- case AtomicRMWInst::UIncWrap:
- Encoding = bitc::RMW_UINC_WRAP;
- break;
- case AtomicRMWInst::UDecWrap:
- Encoding = bitc::RMW_UDEC_WRAP;
- break;
- case AtomicRMWInst::USubCond:
- Encoding = bitc::RMW_USUB_COND;
- break;
- case AtomicRMWInst::USubSat:
- Encoding = bitc::RMW_USUB_SAT;
- break;
- }
-
- if (I.isElementwise())
- Encoding |= bitc::RMW_ELEMENTWISE_FLAG;
- return Encoding;
-}
-
static unsigned getEncodedOrdering(AtomicOrdering Ordering) {
switch (Ordering) {
case AtomicOrdering::NotAtomic: return bitc::ORDERING_NOTATOMIC;
@@ -1860,6 +1780,88 @@ static uint64_t getOptimizationFlags(const Value *V) {
return Flags;
}
+static unsigned getEncodedRMWOperation(const AtomicRMWInst &I) {
+ unsigned Encoding = 0;
+ switch (I.getOperation()) {
+ default: llvm_unreachable("Unknown RMW operation!");
+ case AtomicRMWInst::Xchg:
+ Encoding = bitc::RMW_XCHG;
+ break;
+ case AtomicRMWInst::Add:
+ Encoding = bitc::RMW_ADD;
+ break;
+ case AtomicRMWInst::Sub:
+ Encoding = bitc::RMW_SUB;
+ break;
+ case AtomicRMWInst::And:
+ Encoding = bitc::RMW_AND;
+ break;
+ case AtomicRMWInst::Nand:
+ Encoding = bitc::RMW_NAND;
+ break;
+ case AtomicRMWInst::Or:
+ Encoding = bitc::RMW_OR;
+ break;
+ case AtomicRMWInst::Xor:
+ Encoding = bitc::RMW_XOR;
+ break;
+ case AtomicRMWInst::Max:
+ Encoding = bitc::RMW_MAX;
+ break;
+ case AtomicRMWInst::Min:
+ Encoding = bitc::RMW_MIN;
+ break;
+ case AtomicRMWInst::UMax:
+ Encoding = bitc::RMW_UMAX;
+ break;
+ case AtomicRMWInst::UMin:
+ Encoding = bitc::RMW_UMIN;
+ break;
+ case AtomicRMWInst::FAdd:
+ Encoding = bitc::RMW_FADD;
+ break;
+ case AtomicRMWInst::FSub:
+ Encoding = bitc::RMW_FSUB;
+ break;
+ case AtomicRMWInst::FMax:
+ Encoding = bitc::RMW_FMAX;
+ break;
+ case AtomicRMWInst::FMin:
+ Encoding = bitc::RMW_FMIN;
+ break;
+ case AtomicRMWInst::FMaximum:
+ Encoding = bitc::RMW_FMAXIMUM;
+ break;
+ case AtomicRMWInst::FMinimum:
+ Encoding = bitc::RMW_FMINIMUM;
+ break;
+ case AtomicRMWInst::FMaximumNum:
+ Encoding = bitc::RMW_FMAXIMUMNUM;
+ break;
+ case AtomicRMWInst::FMinimumNum:
+ Encoding = bitc::RMW_FMINIMUMNUM;
+ break;
+ case AtomicRMWInst::UIncWrap:
+ Encoding = bitc::RMW_UINC_WRAP;
+ break;
+ case AtomicRMWInst::UDecWrap:
+ Encoding = bitc::RMW_UDEC_WRAP;
+ break;
+ case AtomicRMWInst::USubCond:
+ Encoding = bitc::RMW_USUB_COND;
+ break;
+ case AtomicRMWInst::USubSat:
+ Encoding = bitc::RMW_USUB_SAT;
+ break;
+ }
+
+ if (I.isElementwise())
+ Encoding |= bitc::RMW_ELEMENTWISE_FLAG;
+ if (getOptimizationFlags(&I) != 0)
+ Encoding |= bitc::RMW_FMF_FLAG;
+ return Encoding;
+}
+
void ModuleBitcodeWriter::writeValueAsMetadata(
const ValueAsMetadata *MD, SmallVectorImpl<uint64_t> &Record) {
// Mimic an MDNode with a value as one operand.
@@ -3624,17 +3626,21 @@ void ModuleBitcodeWriter::writeInstruction(const Instruction &I,
Vals.push_back(cast<AtomicCmpXchgInst>(I).isWeak());
Vals.push_back(getEncodedAlign(cast<AtomicCmpXchgInst>(I).getAlign()));
break;
- case Instruction::AtomicRMW:
+ case Instruction::AtomicRMW: {
+ const AtomicRMWInst &RMWI = cast<AtomicRMWInst>(I);
Code = bitc::FUNC_CODE_INST_ATOMICRMW;
pushValueAndType(I.getOperand(0), InstID, Vals); // ptrty + ptr
pushValueAndType(I.getOperand(1), InstID, Vals); // valty + val
- Vals.push_back(getEncodedRMWOperation(cast<AtomicRMWInst>(I)));
- Vals.push_back(cast<AtomicRMWInst>(I).isVolatile());
- Vals.push_back(getEncodedOrdering(cast<AtomicRMWInst>(I).getOrdering()));
- Vals.push_back(
- getEncodedSyncScopeID(cast<AtomicRMWInst>(I).getSyncScopeID()));
- Vals.push_back(getEncodedAlign(cast<AtomicRMWInst>(I).getAlign()));
+ uint64_t Flags = getOptimizationFlags(&I);
+ Vals.push_back(getEncodedRMWOperation(RMWI));
+ if (Flags != 0)
+ Vals.push_back(Flags);
+ Vals.push_back(RMWI.isVolatile());
+ Vals.push_back(getEncodedOrdering(RMWI.getOrdering()));
+ Vals.push_back(getEncodedSyncScopeID(RMWI.getSyncScopeID()));
+ Vals.push_back(getEncodedAlign(RMWI.getAlign()));
break;
+ }
case Instruction::Fence:
Code = bitc::FUNC_CODE_INST_FENCE;
Vals.push_back(getEncodedOrdering(cast<FenceInst>(I).getOrdering()));
diff --git a/llvm/lib/CodeGen/AtomicExpandPass.cpp b/llvm/lib/CodeGen/AtomicExpandPass.cpp
index 8f75462250f82..b22faed874926 100644
--- a/llvm/lib/CodeGen/AtomicExpandPass.cpp
+++ b/llvm/lib/CodeGen/AtomicExpandPass.cpp
@@ -789,7 +789,8 @@ bool AtomicExpandImpl::tryExpandAtomicRMW(AtomicRMWInst *AI) {
} else {
auto PerformOp = [&](IRBuilderBase &Builder, Value *Loaded) {
return buildAtomicRMWValue(AI->getOperation(), Builder, Loaded,
- AI->getValOperand());
+ AI->getValOperand(),
+ AI->getFastMathFlagsOrNone());
};
expandAtomicOpToLLSC(AI, AI->getType(), AI->getPointerOperand(),
AI->getAlign(), AI->getOrdering(), PerformOp);
@@ -1013,7 +1014,8 @@ static Value *insertMaskedValue(IRBuilderBase &Builder, Value *WideWord,
static Value *performMaskedAtomicOp(AtomicRMWInst::BinOp Op,
IRBuilderBase &Builder, Value *Loaded,
Value *Shifted_Inc, Value *Inc,
- const PartwordMaskValues &PMV) {
+ const PartwordMaskValues &PMV,
+ FastMathFlags FMF) {
// TODO: update to use
// https://graphics.stanford.edu/~seander/bithacks.html#MaskedMerge in order
// to merge bits from two values without requiring PMV.Inv_Mask.
@@ -1031,7 +1033,8 @@ static Value *performMaskedAtomicOp(AtomicRMWInst::BinOp Op,
case AtomicRMWInst::Sub:
case AtomicRMWInst::Nand: {
// The other arithmetic ops need to be masked into place.
- Value *NewVal = buildAtomicRMWValue(Op, Builder, Loaded, Shifted_Inc);
+ Value *NewVal =
+ buildAtomicRMWValue(Op, Builder, Loaded, Shifted_Inc, FMF);
Value *NewVal_Masked = Builder.CreateAnd(NewVal, PMV.Mask);
Value *Loaded_MaskOut = Builder.CreateAnd(Loaded, PMV.Inv_Mask);
Value *FinalVal = Builder.CreateOr(Loaded_MaskOut, NewVal_Masked);
@@ -1057,7 +1060,8 @@ static Value *performMaskedAtomicOp(AtomicRMWInst::BinOp Op,
// the original size, and expand out again after doing the
// operation. Bitcasts will be inserted for FP values.
Value *Loaded_Extract = extractMaskedValue(Builder, Loaded, PMV);
- Value *NewVal = buildAtomicRMWValue(Op, Builder, Loaded_Extract, Inc);
+ Value *NewVal =
+ buildAtomicRMWValue(Op, Builder, Loaded_Extract, Inc, FMF);
Value *FinalVal = insertMaskedValue(Builder, Loaded, NewVal, PMV);
return FinalVal;
}
@@ -1102,7 +1106,8 @@ void AtomicExpandImpl::expandPartwordAtomicRMW(
auto PerformPartwordOp = [&](IRBuilderBase &Builder, Value *Loaded) {
return performMaskedAtomicOp(Op, Builder, Loaded, ValOperand_Shifted,
- AI->getValOperand(), PMV);
+ AI->getValOperand(), PMV,
+ AI->getFastMathFlagsOrNone());
};
Value *OldResult;
@@ -1865,7 +1870,8 @@ bool AtomicExpandImpl::expandAtomicRMWToCmpXchg(
AI->getOrdering(), AI->getSyncScopeID(), AI->isVolatile(),
[&](IRBuilderBase &Builder, Value *Loaded) {
return buildAtomicRMWValue(AI->getOperation(), Builder, Loaded,
- AI->getValOperand());
+ AI->getValOperand(),
+ AI->getFastMathFlagsOrNone());
},
CreateCmpXchg, /*MetadataSrc=*/AI);
diff --git a/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp b/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp
index c66226258c87e..97bf3f1079956 100644
--- a/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp
+++ b/llvm/lib/CodeGen/GlobalISel/IRTranslator.cpp
@@ -3684,7 +3684,8 @@ bool IRTranslator::translateAtomicRMW(const User &U,
*MF->getMachineMemOperand(MachinePointerInfo(I.getPointerOperand()),
Flags, MRI->getType(Val), getMemOpAlign(I),
I.getAAMetadata(), nullptr, I.getSyncScopeID(),
- I.getOrdering()));
+ I.getOrdering()),
+ MachineInstr::copyFlagsFromInstruction(I));
return true;
}
diff --git a/llvm/lib/CodeGen/GlobalISel/MachineIRBuilder.cpp b/llvm/lib/CodeGen/GlobalISel/MachineIRBuilder.cpp
index bd6cd3c325951..2bf1864faeacf 100644
--- a/llvm/lib/CodeGen/GlobalISel/MachineIRBuilder.cpp
+++ b/llvm/lib/CodeGen/GlobalISel/MachineIRBuilder.cpp
@@ -1078,7 +1078,7 @@ MachineIRBuilder::buildAtomicCmpXchg(const DstOp &OldValRes, const SrcOp &Addr,
MachineInstrBuilder MachineIRBuilder::buildAtomicRMW(
unsigned Opcode, const DstOp &OldValRes,
const SrcOp &Addr, const SrcOp &Val,
- MachineMemOperand &MMO) {
+ MachineMemOperand &MMO, std::optional<unsigned> Flags) {
#ifndef NDEBUG
LLT OldValResTy = OldValRes.getLLTTy(*getMRI());
@@ -1095,6 +1095,8 @@ MachineInstrBuilder MachineIRBuilder::buildAtomicRMW(
Addr.addSrcToMIB(MIB);
Val.addSrcToMIB(MIB);
MIB.addMemOperand(&MMO);
+ if (Flags)
+ MIB->setFlags(*Flags);
return MIB;
}
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
index c61a2edd45255..3134093df01d7 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp
@@ -959,6 +959,16 @@ static void AddNodeIDCustom(FoldingSetNodeID &ID, const SDNode *N) {
case ISD::ATOMIC_LOAD_MAX:
case ISD::ATOMIC_LOAD_UMIN:
case ISD::ATOMIC_LOAD_UMAX:
+ case ISD::ATOMIC_LOAD_FADD:
+ case ISD::ATOMIC_LOAD_FSUB:
+ case ISD::ATOMIC_LOAD_FMAX:
+ case ISD::ATOMIC_LOAD_FMIN:
+ case ISD::ATOMIC_LOAD_FMAXIMUM:
+ case ISD::ATOMIC_LOAD_FMINIMUM:
+ case ISD::ATOMIC_LOAD_UINC_WRAP:
+ case ISD::ATOMIC_LOAD_UDEC_WRAP:
+ case ISD::ATOMIC_LOAD_USUB_COND:
+ case ISD::ATOMIC_LOAD_USUB_SAT:
case ISD::ATOMIC_LOAD:
case ISD::ATOMIC_STORE: {
const AtomicSDNode *AT = cast<AtomicSDNode>(N);
@@ -10378,7 +10388,8 @@ SDValue SelectionDAG::getAtomicMemset(SDValue Chain, const SDLoc &dl,
SDValue SelectionDAG::getAtomic(unsigned Opcode, const SDLoc &dl, EVT MemVT,
SDVTList VTList, ArrayRef<SDValue> Ops,
MachineMemOperand *MMO,
- ISD::LoadExtType ExtType) {
+ ISD::LoadExtType ExtType,
+ const SDNodeFlags Flags) {
FoldingSetNodeID ID;
AddNodeIDNode(ID, Opcode, VTList, Ops);
ID.AddInteger(MemVT.getRawBits());
@@ -10390,12 +10401,14 @@ SDValue SelectionDAG::getAtomic(unsigned Opcode, const SDLoc &dl, EVT MemVT,
if (auto *E = cast_or_null<AtomicSDNode>(FindNodeOrInsertPos(ID, dl, IP))) {
E->refineAlignment(MMO);
E->refineRanges(MMO);
+ E->intersectFlagsWith(Flags);
return SDValue(E, 0);
}
auto *N = newSDNode<AtomicSDNode>(dl.getIROrder(), dl.getDebugLoc(), Opcode,
VTList, MemVT, MMO, ExtType);
createOperands(N, Ops);
+ N->setFlags(Flags);
CSEMap.InsertNode(N, IP);
InsertNode(N);
@@ -10418,7 +10431,8 @@ SDValue SelectionDAG::getAtomicCmpSwap(unsigned Opcode, const SDLoc &dl,
SDValue SelectionDAG::getAtomic(unsigned Opcode, const SDLoc &dl, EVT MemVT,
SDValue Chain, SDValue Ptr, SDValue Val,
- MachineMemOperand *MMO) {
+ MachineMemOperand *MMO,
+ const SDNodeFlags Flags) {
assert((Opcode == ISD::ATOMIC_LOAD_ADD || Opcode == ISD::ATOMIC_LOAD_SUB ||
Opcode == ISD::ATOMIC_LOAD_AND || Opcode == ISD::ATOMIC_LOAD_CLR ||
Opcode == ISD::ATOMIC_LOAD_OR || Opcode == ISD::ATOMIC_LOAD_XOR ||
@@ -10441,7 +10455,7 @@ SDValue SelectionDAG::getAtomic(unsigned Opcode, const SDLoc &dl, EVT MemVT,
SDVTList VTs = Opcode == ISD::ATOMIC_STORE ? getVTList(MVT::Other) :
getVTList(VT, MVT::Other);
SDValue Ops[] = {Chain, Ptr, Val};
- return getAtomic(Opcode, dl, MemVT, VTs, Ops, MMO);
+ return getAtomic(Opcode, dl, MemVT, VTs, Ops, MMO, ISD::NON_EXTLOAD, Flags);
}
SDValue SelectionDAG::getAtomicLoad(ISD::LoadExtType ExtType, const SDLoc &dl,
diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
index 0f6bd53cafdd8..117cbb39d99fe 100644
--- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
@@ -5335,10 +5335,13 @@ void SelectionDAGBuilder::visitAtomicRMW(const AtomicRMWInst &I) {
MachinePointerInfo(I.getPointerOperand()), Flags, MemVT.getStoreSize(),
I.getAlign(), AAMDNodes(), nullptr, SSID, Ordering);
- SDValue L =
- DAG.getAtomic(NT, dl, MemVT, InChain,
- getValue(I.getPointerOperand()), getValue(I.getValOperand()),
- MMO);
+ SDNodeFlags NodeFlags;
+ if (auto *FPMO = dyn_cast<FPMathOperator>(&I))
+ NodeFlags.copyFMF(*FPMO);
+
+ SDValue L = DAG.getAtomic(NT, dl, MemVT, InChain,
+ getValue(I.getPointerOperand()),
+ getValue(I.getValOperand()), MMO, NodeFlags);
SDValue OutChain = L.getValue(1);
diff --git a/llvm/lib/IR/Instructions.cpp b/llvm/lib/IR/Instructions.cpp
index e40aae61c723c..1f9a5807190df 100644
--- a/llvm/lib/IR/Instructions.cpp
+++ b/llvm/lib/IR/Instructions.cpp
@@ -4478,6 +4478,7 @@ AtomicRMWInst *AtomicRMWInst::cloneImpl() const {
getOperation(), getOperand(0), getOperand(1), getAlign(), getOrdering(),
getSyncScopeID(), isElementwise());
Result->setVolatile(isVolatile());
+ Result->FMF = FMF;
return Result;
}
diff --git a/llvm/lib/IR/Operator.cpp b/llvm/lib/IR/Operator.cpp
index 6235ef7000c73..eba379b0ade0a 100644
--- a/llvm/lib/IR/Operator.cpp
+++ b/llvm/lib/IR/Operator.cpp
@@ -321,6 +321,8 @@ FastMathFlags &FPMathOperator::getFastMathFlagsImpl() {
return Op->FMF;
if (FastMathFlagsStorage *Op = dyn_cast<CallInst>(I))
return Op->FMF;
+ if (FastMathFlagsStorage *Op = dyn_cast<AtomicRMWInst>(I))
+ return Op->FMF;
if (FastMathFlagsStorage *Op = dyn_cast<UIToFPInst>(I))
return Op->FMF;
if (FastMathFlagsStorage *Op = dyn_cast<SIToFPInst>(I))
diff --git a/llvm/lib/Transforms/Utils/LowerAtomic.cpp b/llvm/lib/Transforms/Utils/LowerAtomic.cpp
index d4c0602f6ce52..27cd6b18ca9ee 100644
--- a/llvm/lib/Transforms/Utils/LowerAtomic.cpp
+++ b/llvm/lib/Transforms/Utils/LowerAtomic.cpp
@@ -51,7 +51,7 @@ std::pair<Value *, Value *> llvm::buildCmpXchgValue(IRBuilderBase &Builder,
Value *llvm::buildAtomicRMWValue(AtomicRMWInst::BinOp Op,
IRBuilderBase &Builder, Value *Loaded,
- Value *Val) {
+ Value *Val, FastMathFlags FMF) {
Value *NewVal;
switch (Op) {
case AtomicRMWInst::Xchg:
@@ -81,21 +81,21 @@ Value *llvm::buildAtomicRMWValue(AtomicRMWInst::BinOp Op,
NewVal = Builder.CreateICmpULE(Loaded, Val);
return Builder.CreateSelect(NewVal, Loaded, Val, "new");
case AtomicRMWInst::FAdd:
- return Builder.CreateFAdd(Loaded, Val, "new");
+ return Builder.CreateFAddFMF(Loaded, Val, FMF, "new");
case AtomicRMWInst::FSub:
- return Builder.CreateFSub(Loaded, Val, "new");
+ return Builder.CreateFSubFMF(Loaded, Val, FMF, "new");
case AtomicRMWInst::FMax:
- return Builder.CreateMaxNum(Loaded, Val);
+ return Builder.CreateMaxNum(Loaded, Val, FMF);
case AtomicRMWInst::FMin:
- return Builder.CreateMinNum(Loaded, Val);
+ return Builder.CreateMinNum(Loaded, Val, FMF);
case AtomicRMWInst::FMaximum:
- return Builder.CreateMaximum(Loaded, Val);
+ return Builder.CreateMaximum(Loaded, Val, FMF);
case AtomicRMWInst::FMinimum:
- return Builder.CreateMinimum(Loaded, Val);
+ return Builder.CreateMinimum(Loaded, Val, FMF);
case AtomicRMWInst::FMaximumNum:
- return Builder.CreateMaximumNum(Loaded, Val);
+ return Builder.CreateMaximumNum(Loaded, Val, FMF);
case AtomicRMWInst::FMinimumNum:
- return Builder.CreateMinimumNum(Loaded, Val);
+ return Builder.CreateMinimumNum(Loaded, Val, FMF);
case AtomicRMWInst::UIncWrap: {
Constant *One = ConstantInt::get(Loaded->getType(), 1);
Value *Inc = Builder.CreateAdd(Loaded, One);
@@ -135,7 +135,8 @@ bool llvm::lowerAtomicRMWInst(AtomicRMWInst *RMWI) {
Value *Val = RMWI->getValOperand();
LoadInst *Orig = Builder.CreateLoad(Val->getType(), Ptr);
- Value *Res = buildAtomicRMWValue(RMWI->getOperation(), Builder, Orig, Val);
+ Value *Res = buildAtomicRMWValue(RMWI->getOperation(), Builder, Orig, Val,
+ RMWI->getFastMathFlagsOrNone());
Builder.CreateStore(Res, Ptr);
RMWI->replaceAllUsesWith(Orig);
RMWI->eraseFromParent();
diff --git a/llvm/test/Assembler/invalid-atomicrmw-elementwise.ll b/llvm/test/Assembler/invalid-atomicrmw.ll
similarity index 81%
rename from llvm/test/Assembler/invalid-atomicrmw-elementwise.ll
rename to llvm/test/Assembler/invalid-atomicrmw.ll
index 2900779420b0c..a668e656f2768 100644
--- a/llvm/test/Assembler/invalid-atomicrmw-elementwise.ll
+++ b/llvm/test/Assembler/invalid-atomicrmw.ll
@@ -3,6 +3,7 @@
; RUN: not llvm-as -disable-output %t/odd-sized.ll 2>&1 | FileCheck %t/odd-sized.ll
; RUN: not llvm-as -disable-output %t/add-must-be-integer.ll 2>&1 | FileCheck %t/add-must-be-integer.ll
; RUN: not llvm-as -disable-output %t/fadd-must-be-fp.ll 2>&1 | FileCheck %t/fadd-must-be-fp.ll
+; RUN: not llvm-as -disable-output %t/fast-math-integer.ll 2>&1 | FileCheck %t/fast-math-integer.ll
;--- scalar.ll
; CHECK: atomicrmw elementwise operand must be a fixed vector type
@@ -31,3 +32,10 @@ define <4 x i32> @bad_fadd(ptr %p, <4 x i32> %v) {
%old = atomicrmw elementwise fadd ptr %p, <4 x i32> %v monotonic
ret <4 x i32> %old
}
+
+;--- fast-math-integer.ll
+; CHECK: fast-math-flags specified for atomicrmw without floating-point type
+define i32 @bad_fast_math(ptr %p, i32 %v) {
+ %old = atomicrmw nnan add ptr %p, i32 %v monotonic
+ ret i32 %old
+}
diff --git a/llvm/test/Bitcode/Inputs/invalid-atomicrmw-fmf-surplus.bc b/llvm/test/Bitcode/Inputs/invalid-atomicrmw-fmf-surplus.bc
new file mode 100644
index 0000000000000000000000000000000000000000..124396ee62f44d22cef8ca957a8c54e8d82de3a2
GIT binary patch
literal 2028
zcmZ`)e at q+q75|z8&H?kiBclVCyQ?qI%F<|ZHA%R`_GEJ{a<@!#(@3M*y4Zv`WJ+3N
zTrhO%Y|iCK>s9V#rPh>9piWaKwUV=;%73se=9udeD`HVKp$sv*{8%#4;E`mCwEVGm
zW|OvQv*dTa)4TV5 at AKoiPrkN(tOcPugwUW#=={b at zxnwO|33fC_Cn{g23sBCO$e<T
z5NfR9Q4QpDz?Y15yz4I1K3`AUq#qgLrYGx+X0G<_!KmrEn!dL>;+?JWf at R1Wq=m~Z
zLqqhtzEAt6P0a`15r^WQ&y&rEJSU?jGk>{es9<dFlwAGO^mT*H7Wddrc~4p1&${|U
zF_YOgO<n$qZK%!sJhlxj4TO?>M}P83n(Yl2X5LBi&rT<5ZATHhWI?C_dO8%SrGMEs
z^(lQ4LULr$>uEUfl$@{Yvpr>O__4t%Cu7lO$a|bWfq&Kd>!u&Pd1rTQvg_49eS7=p
zPrv)#doN%3xWC~`)~`Xze|$_=qpp8{U_`!|$7*#X{`GMobg;d)VU$E at X%NwXRktz9
zER~r+i@`Y~i*YT{7bPbm<Y;=tw2i1QXna44yG?!EVvfBr#{`7BYMRnsB%B4=nIN3S
zNFp30UpJBy9yJ&wCxgUTL=9S0bFmrkBym}wUJ<B`3a)r at H@(EJ^UNyGF8aBEaLK-)
z*jHp}o^UKEjyPd2R$9h<<m<y^&#)ST7*&HwHRvP9tmN3kIw(6HRZkn$kv27yRwvT2
zuDca{w}c;D!#e^oSeRor=a^zOuKAf|o>_{q%l`U5BxR`}OX+INYGO(jX9#;rI|7G3
znzDaJ*fW~xlk;kDKn(?z at Gu$iso^#?Jgb^-yhz<e)See_Ch>-gDqHXlOkuIIYdLPw
z&E4dgwHUWnVm8a{UkqHJ?FBI-_g+!#MQ`Un2l1)n5i*i~;l%wUr6a(7k6ZRLw|RCw
z$G&0U^oDu1;;%pTrXm(Kdq#EwAdeOuo3ecdZnebk8emd&9QrYmlU6d6j(ul?#yi*W
zmKWa>I&wGN+-A9Twal(V10J(Iqd4c4_P8b{WT~h*vN@@s!HkZgXMV&;LRTXI;DBoW
zV;=ALs4q(O6$`Z`P+J!2Z(jTgjW+}gxA#i*cS?BEiysOdpYQmY#TdI<Zr$P8ZJu4?
zSzz!s|3tv at 6Z?Yf%o0*ck%|Kgz&;E+34MCh at t|tH{%5=^P~{T7<D%}&!YP`882x~R
zSx%=xDmud96`(Q}cS|J=pmPFKMc2|SyFAC1+}wJO33Q$ja9QV<f+Aaj%qy?8k)s~<
z^sM^*D7<?$WpE at kM@kb{<X(_d5fC$Z_{aME9JG2P$DWL}1Wue2 at UDg0E!FROsrz~<
zyj7~-wov6281{XaGB!+37*+Gh$FrIpboet_ECPmvlvEr=!U15WG*f?59;6G3nq%IG
zaeD^fF at Lol6bN;Ii|&KR+tgr`ocw^CumWYddAFD$><gOssK&XV^`@Y=;WZ0;%fJD8
zc>pliq_gjz-m4Zt)W$X4EcXN at rc{~IX9WF^TFJ at uUzgmyS!hhrbvrBfg8M+dn_A-Q
zcsa%v%dOjUOg_flEVFCnRvkKM&&`3%mceK(qe1*2i8qCfIc3bMjs>B<I^t4859`7n
zavXI25(|!iRzFnQKU3fo35w6ly({upvjzv~CoT78fEw?PZ(j+_CeM@%tiDx#(T32s
z;M(|B5Nsoo;{&b)2$zt$BLBybx#uvbZOd1UCA{Iq`@Z^Z1 at BcF?|LZ^cVj{?715!r
zQ}s2Gm7RcQm$(K_%dt?o=r<gtwh(xdDuXgAjrUv>a9j at Jt<0=N4T0-s;av&lNr2l4
zXV>rzKXWU}LF;RN?&dyNes?$C26Sk==fd|P!s{-ou`QBRL%LgZ;{|8ckVUQgL`ul*
zVDD8~>dNF}tiBL)O|S*>5MzKflQ*!-F=k&SUuG&UYLCVbT-5DMbU+;u)bKzBkjW~}
zl;XG|OZp}P)hdpRHg!qgRyu#MR`wg4Yue_YxaDTnV}SDSX at u$lW3uP~1zf3)GpjgY
zktH=);mB`A=m1<hj8j)hmF|ktrl*a#77_5QUl6bR2>02VoPPITxii&V-^x_U^}ljQ
zET~$;hUr6blqcF~XSK%5 at ITuMzYY?qkMswDbGD<0g~P(R^I^-N?U}%!)ZG(!=J+vD
z63_J5dW3GF=eQ6M{GxyGxOnVrKoHLk9tmFz9`T1?y?Fji=;HX(f%B2 at YoW7;FI=de
HfY3hxA7SSn
literal 0
HcmV?d00001
diff --git a/llvm/test/Bitcode/Inputs/invalid-atomicrmw-fmf-truncated.bc b/llvm/test/Bitcode/Inputs/invalid-atomicrmw-fmf-truncated.bc
new file mode 100644
index 0000000000000000000000000000000000000000..3ee403a6e27bb0a2a4950cb19c449c4b714bba5e
GIT binary patch
literal 2028
zcmZ`)e at q+q75|#UoCD^2M at 9!OcUNDam8H?-YLak=?aAg^<ZhYfrjbUqb+HL?$dt6k
zxM1ki*__Lf&a2$XO3{=}MxCZk>LjN{mH%K{%rTb|D`HVKp$sv*{8%#4;E`mCwEVGm
zW|OvQv*dTa)4TV5 at AKoiSH8A>tQjFALTJD&bbRBb-~8f7f1mqi+fv7~2D=gQW`tG^
z2sPC3s0MP{;Y-Hb-|;NfK3_-Mr5_s-<|mCN3s?KYftdNZn%=kC6CEvyCF`IqL<^Uj
z2M6hQ{GawtnVSy0Ee<BUpQoA*c~8d77XEV0;F77SLvr^`(bo-jd%|l!<vV44FX!$H
z$ITY|6m|J4_Q6)m^VmMP*dI>u?R}{yX|^Y_H2rppe|9QaYd?z6B`ZSp(9 at w{E&XBd
z<frsW2&v%(pSS+NQ*y!BYk$gA{}Y2vPQ_zQkoP!$0)N!<>&73vd1rTYqVv^1eS7=p
z&%XQayDwk(xUc?8wy#0Te|*eWWA1-{U_$=s$7*#X{`GMqbg-?qeuPA5aRAYPRhKEo
zES8yIv%xhjiwP~+8zaY~<jBmhc^gq*(D;4~cbR*+g*<y>mI(^RYMRnkBwS0fD at nMD
z(PSh<zHTDNy=o{#PK1cjs2Z}WmSPj$N#U|Uy&_N>6<qP*E_#t&=b2TWT?lYN;gVxs
zajeMF9O0Z-oC(5FtTd1M$=8R-?jbb{F{Xx6YRFHH+Q`v|MkqTLQ%{@J;Z`*~qmIwS
zJMUKT-4cFq4etoVz|t(UIm;BQaV@|s^UPwLT at KX!Atg&ovNThTIZaH<;xyq%Ye(SF
zN0W}v2uD^ke{x<8^{e5K5*Z?+el^mnMl!19#*5TlMD6+TW(se(sj?OCz!X*+yO!q`
zJlsv5S&MUPC1$hC{>8urTVD{fa?cgTQS^2Ea}d8e7A2!IFPylaqI3kf?{Uik<~Gl+
z=h-(5oZc|cRswaW-c-b*=E%w}0OZkvb5nLq!>yM5T|G>yjzK>ra>7Q2XX4-4pz+Q%
zyye69g!cSR54TxvSuL~c(16$C$SSTmr7fX}Nm(jt&Rkwv(qKks(K|P6BB85c0I*-R
z{IP&{{L~kvx{8(B5~wXJ^*0}Wg~l5KhTD6kx;rJj>BA3&_Rn_$%tD-9Ew}9O>^9FX
z at +>fTn|~r`{i$PKcI5~utw_cGd0-!goq#^Q>R3p%T>mrP6{vCv-*HoSGH{A!AV%IN
zVV2V~AQc_q$O=#yPk5w~2GF^HsiJ!^!!FOVB at efrXM!DP1YFknrJ%@`AoI#=t>lPT
zJ)Kd%AA@(VrVY-d=1gniirfQoDgt6=FaKC?fP+?V<k^$)=HQ8Q0^YS!yQR8aA9Y_Z
zg||v|+g7T)0>i%NRz`=&ag%B}`FKuqfDV5yi$%bYkWz}XNH_t^v}W#WEP!-DQM1e&
zac<85JQl9jfdZipaM5+}c&i$UkrVHe<2IlyKj#s%gkxS4AJw?#wVpKeHne7CZy7j1
zuK)n%8+G;rQ+w4Sh}yWOo8_LM#FZ*j`i!9e5gR!%^VdaBPYxPWblqm;9&jJ1cT-D#
z9WTe(V!35|mMO%!n`L&b+ at eDV?Rhwm*)kZdc_f4%r0}M&F{_N))X@;sSBKqd_ at Oc4
zCC5PLFR|bVX!Qf7?K1^Vk)ZgT+_NHoHD_>werDvJEKuXy@$V~v+2om$fz`LlFIy4%
z7F-+N3W04za;)E-1mO}=XY~IVGXER~wQc#Uv4l5#c;8pQt>C>%!(AT*;%-RlrJ_2N
zMpa)EIoSnhc8Y7 at v^)!yivhz?Y72oUsWK>|(s0jB0mtPK-pXdIY8YIXfp;aCCkbvR
zT%AKV0?e%#2d%FKxSRW61w37N8_=Qgo*UnX2(P=ThSq3G4eM^vjTg$OVXJEVL`uqS
zVDD8~>dY46tiBNQjj#m at 5aWO~Q!ud0ab{m7UuG(9YLCVb+|=!CtX~}#)JT5?kjW{o
zwBo!XOZp}P)hf=cHhD?kRyu#MR`wg4ZQSOcxaDEi<ACz-X at u$kW3uQ31zf3)lTn<o
z$Wj`ta1^j1bO5gHrpc at 16#c;@GCgg=wTOUceS&!1Pk7GO<n_D%%AKy}=2s?5?*ElL
zY(>=?c1#~ipaRiKyQ(!_hX2_X_;rv-|4M%lJZC?8SU4=4J0GzQ*q;dwNL}5*XO15e
zCGkwRy<6xKx{nJ%As_|@j*G|61_kl#z>&zs(2+po)r;rPgfEUg9Xub6ycRxt_`-$i
H2?+fI?Z at a^
literal 0
HcmV?d00001
diff --git a/llvm/test/Bitcode/atomic.ll b/llvm/test/Bitcode/atomic.ll
index 3c9ddbd6446ee..9f945f8d8aea6 100644
--- a/llvm/test/Bitcode/atomic.ll
+++ b/llvm/test/Bitcode/atomic.ll
@@ -16,3 +16,22 @@ define void @test_cmpxchg(i32* %addr, i32 %desired, i32 %new) {
ret void
}
+
+define float @test_atomicrmw_fmf(ptr %addr, float %value) {
+ ; CHECK: %fast = atomicrmw fast fadd ptr %addr, float %value monotonic
+ %fast = atomicrmw fast fadd ptr %addr, float %value monotonic
+
+ ; CHECK: %flags = atomicrmw volatile nnan ninf fsub ptr %addr, float %value acquire, align 4
+ %flags = atomicrmw volatile ninf nnan fsub ptr %addr, float %value acquire, align 4
+
+ ; CHECK: %xchg = atomicrmw nsz xchg ptr %addr, float %value seq_cst
+ %xchg = atomicrmw nsz xchg ptr %addr, float %value seq_cst
+ ret float %xchg
+}
+
+define <2 x half> @test_atomicrmw_elementwise_fmf(ptr %addr,
+ <2 x half> %value) {
+ ; CHECK: %old = atomicrmw volatile fast elementwise fadd ptr %addr, <2 x half> %value monotonic
+ %old = atomicrmw volatile fast elementwise fadd ptr %addr, <2 x half> %value monotonic
+ ret <2 x half> %old
+}
diff --git a/llvm/test/Bitcode/invalid.test b/llvm/test/Bitcode/invalid.test
index 718161d20c38f..8cefcf58cd7af 100644
--- a/llvm/test/Bitcode/invalid.test
+++ b/llvm/test/Bitcode/invalid.test
@@ -225,6 +225,13 @@ RUN: FileCheck --check-prefix=NONPOINTER-ATOMICRMW %s
NONPOINTER-ATOMICRMW: Invalid atomicrmw record
+RUN: not llvm-dis -disable-output %p/Inputs/invalid-atomicrmw-fmf-truncated.bc 2>&1 | \
+RUN: FileCheck --check-prefix=INVALID-ATOMICRMW-FMF %s
+RUN: not llvm-dis -disable-output %p/Inputs/invalid-atomicrmw-fmf-surplus.bc 2>&1 | \
+RUN: FileCheck --check-prefix=INVALID-ATOMICRMW-FMF %s
+
+INVALID-ATOMICRMW-FMF: Invalid atomicrmw record
+
RUN: not llvm-dis -disable-output %p/Inputs/invalid-fcmp-opnum.bc 2>&1 | \
RUN: FileCheck --check-prefix=INVALID-FCMP-OPNUM %s
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll
index d0be9b9c66f9f..7d248e44cb369 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-atomicrmw.ll
@@ -8,10 +8,10 @@ define float @test_atomicrmw_fadd(ptr addrspace(3) %addr) {
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[COPY:%[0-9]+]]:_(p3) = COPY $vgpr0
; CHECK-NEXT: [[C:%[0-9]+]]:_(f32) = G_FCONSTANT float 1.000000e+00
- ; CHECK-NEXT: [[ATOMICRMW_FADD:%[0-9]+]]:_(f32) = G_ATOMICRMW_FADD [[COPY]](p3), [[C]] :: (load store seq_cst (f32) on %ir.addr, addrspace 3)
+ ; CHECK-NEXT: [[ATOMICRMW_FADD:%[0-9]+]]:_(f32) = nnan nsz G_ATOMICRMW_FADD [[COPY]](p3), [[C]] :: (load store seq_cst (f32) on %ir.addr, addrspace 3)
; CHECK-NEXT: $vgpr0 = COPY [[ATOMICRMW_FADD]](f32)
; CHECK-NEXT: SI_RETURN implicit $vgpr0
- %oldval = atomicrmw fadd ptr addrspace(3) %addr, float 1.0 seq_cst
+ %oldval = atomicrmw nnan nsz fadd ptr addrspace(3) %addr, float 1.0 seq_cst
ret float %oldval
}
diff --git a/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll b/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll
index 66f1a9b5439f2..22d31ba61c008 100644
--- a/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll
+++ b/llvm/test/Transforms/AtomicExpand/AArch64/atomicrmw-fp.ll
@@ -7,7 +7,7 @@ define float @test_atomicrmw_fadd_f32(ptr %ptr, float %value) !prof !0 {
; CHECK-NEXT: br label [[ATOMICRMW_START:%.*]]
; CHECK: atomicrmw.start:
; CHECK-NEXT: [[LOADED:%.*]] = phi float [ [[TMP1]], [[TMP0:%.*]] ], [ [[TMP5:%.*]], [[ATOMICRMW_START]] ]
-; CHECK-NEXT: [[NEW:%.*]] = fadd float [[LOADED]], [[VALUE:%.*]]
+; CHECK-NEXT: [[NEW:%.*]] = fadd nnan nsz float [[LOADED]], [[VALUE:%.*]]
; CHECK-NEXT: [[TMP2:%.*]] = bitcast float [[NEW]] to i32
; CHECK-NEXT: [[TMP3:%.*]] = bitcast float [[LOADED]] to i32
; CHECK-NEXT: [[TMP4:%.*]] = cmpxchg ptr [[PTR]], i32 [[TMP3]], i32 [[TMP2]] seq_cst seq_cst, align 4
@@ -18,7 +18,7 @@ define float @test_atomicrmw_fadd_f32(ptr %ptr, float %value) !prof !0 {
; CHECK: atomicrmw.end:
; CHECK-NEXT: ret float [[TMP5]]
;
- %res = atomicrmw fadd ptr %ptr, float %value seq_cst
+ %res = atomicrmw nnan nsz fadd ptr %ptr, float %value seq_cst
ret float %res
}
diff --git a/llvm/test/Transforms/LowerAtomic/atomic-load.ll b/llvm/test/Transforms/LowerAtomic/atomic-load.ll
index cf41d14547d2e..837d7f28f5bb5 100644
--- a/llvm/test/Transforms/LowerAtomic/atomic-load.ll
+++ b/llvm/test/Transforms/LowerAtomic/atomic-load.ll
@@ -38,9 +38,9 @@ define i8 @min() {
define float @fadd() {
; CHECK-LABEL: @fadd(
%i = alloca float
- %j = atomicrmw fadd ptr %i, float 42.0 monotonic
+ %j = atomicrmw nnan nsz fadd ptr %i, float 42.0 monotonic
; CHECK: [[INST:%[a-z0-9]+]] = load
-; CHECK-NEXT: fadd
+; CHECK-NEXT: fadd nnan nsz
; CHECK-NEXT: store
ret float %j
; CHECK: ret float [[INST]]
@@ -60,9 +60,9 @@ define float @fsub() {
define float @fmax() {
; CHECK-LABEL: @fmax(
%i = alloca float
- %j = atomicrmw fmax ptr %i, float 42.0 monotonic
+ %j = atomicrmw ninf fmax ptr %i, float 42.0 monotonic
; CHECK: [[INST:%[a-z0-9]+]] = load
-; CHECK-NEXT: call float @llvm.maxnum.f32
+; CHECK-NEXT: call ninf float @llvm.maxnum.f32
; CHECK-NEXT: store
ret float %j
; CHECK: ret float [[INST]]
@@ -82,9 +82,9 @@ define float @fmin() {
define float @fmaximum() {
; CHECK-LABEL: @fmaximum(
%i = alloca float
- %j = atomicrmw fmaximum ptr %i, float 42.0 monotonic
+ %j = atomicrmw nnan fmaximum ptr %i, float 42.0 monotonic
; CHECK: [[INST:%[a-z0-9]+]] = load
-; CHECK-NEXT: call float @llvm.maximum.f32
+; CHECK-NEXT: call nnan float @llvm.maximum.f32
; CHECK-NEXT: store
ret float %j
; CHECK: ret float [[INST]]
diff --git a/llvm/unittests/CodeGen/SelectionDAGNodeConstructionTest.cpp b/llvm/unittests/CodeGen/SelectionDAGNodeConstructionTest.cpp
index 0899b04bfddb8..ad9d6c7304b20 100644
--- a/llvm/unittests/CodeGen/SelectionDAGNodeConstructionTest.cpp
+++ b/llvm/unittests/CodeGen/SelectionDAGNodeConstructionTest.cpp
@@ -47,6 +47,33 @@ TEST_F(SelectionDAGNodeConstructionTest, AND) {
EXPECT_EQ(DAG->getNode(ISD::AND, DL, MVT::i32, Undef, Undef), Undef);
}
+TEST_F(SelectionDAGNodeConstructionTest, AtomicRMWFlags) {
+ SDLoc DL;
+ SDValue Ptr = DAG->getUNDEF(MVT::i64);
+ SDValue Val = DAG->getUNDEF(MVT::f32);
+ MachineMemOperand *MMO = MF->getMachineMemOperand(
+ MachinePointerInfo(),
+ MachineMemOperand::MOLoad | MachineMemOperand::MOStore, 4, Align(4),
+ AAMDNodes(), nullptr, SyncScope::System, AtomicOrdering::Monotonic);
+
+ SDNodeFlags Flags;
+ Flags.setNoNaNs(true);
+ Flags.setNoSignedZeros(true);
+ SDValue Atomic = DAG->getAtomic(ISD::ATOMIC_LOAD_FADD, DL, MVT::f32,
+ DAG->getEntryNode(), Ptr, Val, MMO, Flags);
+ EXPECT_TRUE(Atomic->getFlags().hasNoNaNs());
+ EXPECT_TRUE(Atomic->getFlags().hasNoSignedZeros());
+
+ SDNodeFlags FewerFlags;
+ FewerFlags.setNoNaNs(true);
+ SDValue SameAtomic = DAG->getAtomic(
+ ISD::ATOMIC_LOAD_FADD, DL, MVT::f32, DAG->getEntryNode(), Ptr, Val, MMO,
+ FewerFlags);
+ EXPECT_EQ(Atomic, SameAtomic);
+ EXPECT_TRUE(SameAtomic->getFlags().hasNoNaNs());
+ EXPECT_FALSE(SameAtomic->getFlags().hasNoSignedZeros());
+}
+
TEST_F(SelectionDAGNodeConstructionTest, MUL) {
SDLoc DL;
SDValue Op = DAG->getCopyFromReg(DAG->getEntryNode(), DL,
diff --git a/llvm/unittests/IR/IRBuilderTest.cpp b/llvm/unittests/IR/IRBuilderTest.cpp
index 6befc403c6f30..e5376489a64e3 100644
--- a/llvm/unittests/IR/IRBuilderTest.cpp
+++ b/llvm/unittests/IR/IRBuilderTest.cpp
@@ -864,6 +864,26 @@ TEST_F(IRBuilderTest, FastMathFlags) {
EXPECT_TRUE(FRem->hasNoNaNs());
}
+TEST_F(IRBuilderTest, AtomicRMWFastMathFlags) {
+ IRBuilder<> Builder(BB);
+ auto *AtomicGV = new GlobalVariable(*M, Type::getFloatTy(Ctx), false,
+ GlobalValue::ExternalLinkage, nullptr);
+ Value *Val = Builder.CreateLoad(AtomicGV->getValueType(), AtomicGV);
+
+ FastMathFlags FMF;
+ FMF.setNoNaNs();
+ FMF.setNoSignedZeros();
+ Builder.setFastMathFlags(FMF);
+
+ AtomicRMWInst *RMW = Builder.CreateAtomicRMW(
+ AtomicRMWInst::FAdd, AtomicGV, Val, Align(4),
+ AtomicOrdering::Monotonic);
+ EXPECT_EQ(FMF, RMW->getFastMathFlags());
+
+ std::unique_ptr<AtomicRMWInst> Clone(cast<AtomicRMWInst>(RMW->clone()));
+ EXPECT_EQ(FMF, Clone->getFastMathFlags());
+}
+
TEST_F(IRBuilderTest, WrapFlags) {
IRBuilder<NoFolder> Builder(BB);
More information about the cfe-commits
mailing list