[clang] Add clang builtins support for gfx90a

[gandhi21299]

Summary

Implement target builtins for gfx90a including fadd64, fadd32, add2h, max and min on various global, flat and ds address spaces for which intrinsics are already implemented.

@rampitec Compiler recommended me to add global-noret target feature after setting it in BuiltinsAMDGPU.def. I am not sure what that means outside of the BuiltinsAMDGPU.def so I have changed it back to gfx90a-insts.

[gandhi21299]

• added more tests with supported combinations of address spaces
• minor nits

[gandhi21299]

• added scope argument
• removed irrelevant tests and updated them
• replaced constant by generic qualifier for two of the test functions

[gandhi21299]

• removed kernel from functions taking in __generic qualified addr

[gandhi21299]

• eliminated scope argument, seemed irrelevant

[gandhi21299]

fixed tests by replacing checks for the functions with generic-qualified arguments.

[gandhi21299]

Passed PSDB

[arsenm]

Missing sema test for rejecting the builtins on targets that don't support this

clang/include/clang/Basic/BuiltinsAMDGPU.def

201 ↗

(On Diff #362804)

"_2f16" looks weird to me. The instruction names call it "pk"

clang/lib/CodeGen/CGBuiltin.cpp

16114 ↗

(On Diff #362804)

Initializing this here is strange, sink down to the double case

clang/test/CodeGenOpenCL/builtins-fp-atomics.cl

113

If you're going to bother testing the ISA, is it worth testing rtn and no rtn versions?

[yaxunl]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201 ↗

(On Diff #362804)

This is to have a consistent postfix naming convention, since the stem part here are the same. the postfix is for the argument type of the builtin function.

[gandhi21299]

clang/lib/CodeGen/CGBuiltin.cpp

16114 ↗

(On Diff #362804)

You mean push it down after line 16119?

clang/test/CodeGenOpenCL/builtins-fp-atomics.cl

113

Sorry, what do you mean by rtn version?

[arsenm]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201 ↗

(On Diff #362804)

Just a plain 2 isn't consistent either. The llvm type naming convention would add a v prefix, but the builtins should probably follow the instructions

clang/lib/CodeGen/CGBuiltin.cpp

16114 ↗

(On Diff #362804)

Yes

clang/test/CodeGenOpenCL/builtins-fp-atomics.cl

113

Most atomics can be optimized if they don't return the in memory value if the value is unused

[gandhi21299]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201 ↗

(On Diff #362804)

Yea, v2f16 looks reasonable.

[gandhi21299]

addressed reviewers' feedback:

• changed a builtin name,
• corrected tests,
• minor formatting nits

[gandhi21299]

refreshing patch

[rampitec]

Needs an IR test, a test for different supported targets, and a negative test for unsupported features.

clang/include/clang/Basic/BuiltinsAMDGPU.def

199 ↗

(On Diff #363196)

Correct attribute for this one in atomic-fadd-insts. In particular it was first added in gfx908 and you would need to test it too.

205 ↗

(On Diff #363196)

Flat address space is 0.

210 ↗

(On Diff #363196)

This is available since gfx8. Attribute gfx8-insts.

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

You do not need any of that code. You can directly map a builtin to intrinsic in the IntrinsicsAMDGPU.td.

clang/test/CodeGenOpenCL/builtins-fp-atomics.cl

113

Certainly yes, because global_atomic_add_f32 did not have return version on gfx908.

[gandhi21299]

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

Sorry, I looked around for several days but I could not figure this out. Is there a concrete example?

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

Every instantiation of GCCBuiltin in the IntrinsicsAMDGPU.td.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

This is not true if the intrinsic requires type mangling. GCCBuiltin is too simple to handle it

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

Yes, but these do not need it. All of these builtins are specific.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

These intrinsics are all mangled based on the FP type

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16212 ↗

(On Diff #363196)

Ah, right. Intrinsics are mangled, builtins not. True. OK, this shall be code then.

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16270 ↗

(On Diff #363196)

Should we map flags since we already have them?

[gandhi21299]

clang/lib/CodeGen/CGBuiltin.cpp

16270 ↗

(On Diff #363196)

Do you mean the memory order flag?

[gandhi21299]

@rampitec how do I handle the following?

builtins-fp-atomics.cl:38:10: error: '__builtin_amdgcn_global_atomic_fadd_f64' needs target feature atomic-fadd-insts
  *rtn = __builtin_amdgcn_global_atomic_fadd_f64(addr, x, memory_order_relaxed);
         ^

[rampitec]

In D106909#2922567, @gandhi21299 wrote:

@rampitec how do I handle the following?

builtins-fp-atomics.cl:38:10: error: '__builtin_amdgcn_global_atomic_fadd_f64' needs target feature atomic-fadd-insts
  *rtn = __builtin_amdgcn_global_atomic_fadd_f64(addr, x, memory_order_relaxed);
         ^

It is f64, it needs gfx90a-insts. atomic-fadd-insts is for global f32.

clang/lib/CodeGen/CGBuiltin.cpp

16270 ↗

(On Diff #363196)

All 3: ordering, scope and volatile.

[gandhi21299]

@rampitec what should I be testing exactly in the IR test?

[rampitec]

In D106909#2923724, @gandhi21299 wrote:

@rampitec what should I be testing exactly in the IR test?

Produced call to the intrinsic. All of these tests there doing that.

[gandhi21299]

clang/lib/CodeGen/CGBuiltin.cpp

16270 ↗

(On Diff #363196)

Following the discussion, what change is required here?

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16270 ↗

(On Diff #363196)

Keep zeroes, drop immediate argument of the builtins.

[gandhi21299]

updating test

[gandhi21299]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201 ↗

(On Diff #364152)

I tried add target feature gfx908-insts for this builtin but the frontend complains that it should have target feature gfx90a-insts.

[rampitec]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201 ↗

(On Diff #364152)

That was for global_atomic_fadd_f32, but as per discussion we are going to use builtin only starting from gfx90a because of the noret problem. Comments in the review are off their positions after multiple patch updates.

210 ↗

(On Diff #364152)

This needs tests with a gfx8 target and a negative test with gfx7.

clang/test/CodeGenOpenCL/builtins-fp-atomics-gfx90a.cl

13 ↗

(On Diff #364152)

Use _f64 or _double in the test name.

31 ↗

(On Diff #364152)

Same here and in other double tests, use a suffix f64 or double.

67 ↗

(On Diff #364152)

'constant' is wrong. It is flat. Here and everywhere.

clang/test/CodeGenOpenCL/builtins-fp-atomics-unsupported-gfx908.cl

7 ↗

(On Diff #364152)

Need to check all other buintins too.

[gandhi21299]

• added more negative tests
• fixed some tests

[rampitec]

clang/test/CodeGenOpenCL/builtins-fp-atomics-unsupported-gfx7.cl

8 ↗

(On Diff #364242)

Add new line.

clang/test/CodeGenOpenCL/unsupported-fadd2f16-gfx908.cl

1 ↗

(On Diff #364242)

Combine all of these gfx908 error tests into a single file. For example like in the builtins-amdgcn-dl-insts-err.cl. It is also better to rename these test filenames to follow the existing pattern: builtins-amdgcn-*-err.cl

[gandhi21299]

• combined tests into a single file
• renamed tests for consistency

[rampitec]

LGTM

[gandhi21299]

Thanks a lot for the review, I will merge this patch in :)

[gandhi21299]

I will merge this patch in as soon as the builds are successful.

clang/test/CodeGenCUDA/fp-atomics-optremarks.cu

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:   -target-cpu gfx90a -Rpass=.* -S -o - 2>&1 | \
+// RUN:   FileCheck %s --check-prefix=GFX90A-CAS
+
+// REQUIRES: amdgpu-registered-target
+
+#include "Inputs/cuda.h"
+#include <stdatomic.h>
+
+// GFX90A-CAS: An FP atomic instruction was expanded into a CAS loop.
+// GFX90A-CAS-LABEL: _Z14atomic_add_casPf
+// GFX90A-CAS:  flat_atomic_cmpswap v0, v[2:3], v[4:5] glc
+// GFX90A-CAS:  s_cbranch_execnz
+__device__ float atomic_add_cas(float *p) {
+  return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
+}

clang/test/CodeGenOpenCL/builtins-fp-atomics.cl

+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN:   %s -S -emit-llvm -o - | FileCheck %s -check-prefix=CHECK
+
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN:   -S -o - %s | FileCheck -check-prefix=GFX90A %s
+
+
+typedef enum memory_order {
+  memory_order_relaxed = __ATOMIC_RELAXED,
+  memory_order_acquire = __ATOMIC_ACQUIRE,
+  memory_order_release = __ATOMIC_RELEASE,
+  memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  memory_order_seq_cst = __ATOMIC_SEQ_CST
+} memory_order;
+
+typedef half __attribute__((ext_vector_type(2))) half2;
+
+// CHECK-LABEL: test_global_add
+// CHECK: tail call double @llvm.amdgcn.global.atomic.fadd.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_global_add
+// GFX90A:  global_atomic_add_f64 v2, v[0:1], s[0:1]
+// GFX90A:  s_endpgm
+kernel void test_global_add(__global double *addr, double x) {
+  __builtin_amdgcn_global_atomic_fadd_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_global_addf
+// CHECK: tail call float @llvm.amdgcn.global.atomic.fadd.f32.p1f32.f32(float addrspace(1)* %{{.*}}, float %{{.*}})
+// GFX90A-LABEL: test_global_addf
+// GFX90A: global_atomic_add_f32 v0, v1, s[0:1]
+// GFX90A: s_endpgm
+kernel void test_global_addf(__global float *addr, float x) {
+  __builtin_amdgcn_global_atomic_fadd_f32(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_global_add2h
+// CHECK: tail call <2 x half> @llvm.amdgcn.global.atomic.fadd.v2f16.p1v2f16.v2f16(<2 x half> addrspace(1)* %{{.*}}, <2 x half> %{{.*}})
+// GFX90A-LABEL: test_global_add2h
+// GFX90A: global_atomic_pk_add_f16 v0, v1, s[0:1]
+// GFX90A: s_endpgm
+kernel void test_global_add2h(__global half2 *addr, half2 x){
+  __builtin_amdgcn_global_atomic_fadd_2f16(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_global_global_min
+// CHECK: tail call double @llvm.amdgcn.global.atomic.fmin.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_global_global_min
+// GFX90A:  global_atomic_min_f64 v2, v[0:1], s[0:1]
+// GFX90A:  s_endpgm
+kernel void test_global_global_min(__global double *addr, double x){
+  __builtin_amdgcn_global_atomic_fmin_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_global_max
+// CHECK: tail call double @llvm.amdgcn.global.atomic.fmax.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_global_max
+// GFX90A:  global_atomic_max_f64 v2, v[0:1], s[0:1]
+// GFX90A:  s_endpgm
+kernel void test_global_max(__global double *addr, double x){
+  __builtin_amdgcn_global_atomic_fmax_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_flat_add_local
+// CHECK: tail call double @llvm.amdgcn.flat.atomic.fadd.f64.p3f64.f64(double addrspace(3)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_flat_add_local
+// GFX90A:  ds_add_f64 v2, v[0:1]
+// GFX90A:  s_endpgm
+kernel void test_flat_add_local(__local double *addr, double x){
+  __builtin_amdgcn_flat_atomic_fadd_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_flat_global_add
+// CHECK: tail call double @llvm.amdgcn.flat.atomic.fadd.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_flat_global_add
+// GFX90A:  global_atomic_add_f64
+// GFX90A:  s_endpgm
+kernel void test_flat_global_add(__global double *addr, double x){
+  __builtin_amdgcn_flat_atomic_fadd_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_flat_min_constant
+// CHECK: tail call double @llvm.amdgcn.flat.atomic.fmin.f64.p4f64.f64(double addrspace(4)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_flat_min_constant
+// GFX90A:  global_atomic_min_f64
+// GFX90A:  s_endpgm
+void test_flat_min_constant(__generic double *addr, double x){
+  __builtin_amdgcn_flat_atomic_fmin_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_flat_global_min
+// CHECK: tail call double @llvm.amdgcn.flat.atomic.fmin.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_flat_global_min
+// GFX90A:  global_atomic_min_f64
+// GFX90A:  s_endpgm
+kernel void test_flat_global_min(__global double *addr, double x){
+  __builtin_amdgcn_flat_atomic_fmin_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_flat_max_constant
+// CHECK: tail call double @llvm.amdgcn.flat.atomic.fmax.f64.p4f64.f64(double addrspace(4)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_max_constant
+// GFX90A:  global_atomic_max_f64 v2, v[0:1], s[0:1]
+// GFX90A:  s_endpgm
+void test_flat_max_constant(__generic double *addr, double x){
+  __builtin_amdgcn_flat_atomic_fmax_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_flat_global_max
+// CHECK: tail call double @llvm.amdgcn.flat.atomic.fmax.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_global_max
+// GFX90A:  global_atomic_max_f64 v2, v[0:1], s[0:1]
+// GFX90A:  s_endpgm
+kernel void test_flat_global_max(__global double *addr, double x){
+  __builtin_amdgcn_flat_atomic_fmax_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_ds_add_local
+// CHECK: tail call double @llvm.amdgcn.ds.fadd.f64(double addrspace(3)* %{{.*}}, double %{{.*}},
+// GFX90A:  test_ds_add_local
+// GFX90A:  ds_add_f64 v2, v[0:1]
+// GFX90A:  s_endpgm
+kernel void test_ds_add_local(__local double *addr, double x){
+  __builtin_amdgcn_ds_atomic_fadd_f64(addr, x, memory_order_relaxed, "workgroup");
+}
+
+// CHECK-LABEL: test_ds_addf_local
+// CHECK: tail call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}},
+// GFX90A-LABEL:  test_ds_addf_local
+// GFX90A:  ds_add_f32 v0, v1
+// GFX90A:  s_endpgm
+kernel void test_ds_addf_local(__local float *addr, float x){
+  __builtin_amdgcn_ds_atomic_fadd_f32(addr, x, memory_order_relaxed, "workgroup");
+}

clang/test/CodeGenOpenCL/fp-atomics-optremarks-gfx90a.cl

+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN:     -Rpass=si-lower -munsafe-fp-atomics %s -S -o - 2>&1 \
+// RUN:     | FileCheck %s -check-prefix=GFX90A-HW
+
+typedef enum memory_order {
+  memory_order_relaxed = __ATOMIC_RELAXED,
+  memory_order_acquire = __ATOMIC_ACQUIRE,
+  memory_order_release = __ATOMIC_RELEASE,
+  memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  memory_order_seq_cst = __ATOMIC_SEQ_CST
+} memory_order;
+
+typedef enum memory_scope {
+  memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+  memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+  memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
+  memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+#if defined(cl_intel_subgroups) || defined(cl_khr_subgroups)
+  memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+#endif
+} memory_scope;
+
+// remark: An unsafe hardware instruction was generated.
+// GFX90A-HW-LABEL: test_atomic_add
+// GFX90A-HW:   global_atomic_add_f64
+float test_atomic_add(global atomic_double *d, double a) {
+  return __opencl_atomic_fetch_add(d, a, memory_order_relaxed, memory_scope_work_group);
+}

llvm/lib/Target/AMDGPU/SIISelLowering.h

 
 namespace AMDGPU {
 struct ImageDimIntrinsicInfo;
 }
 
 class SITargetLowering final : public AMDGPUTargetLowering {
 private:
   const GCNSubtarget *Subtarget;
+  OptimizationRemarkEmitter *ORE;
 
 public:
   MVT getRegisterTypeForCallingConv(LLVMContext &Context,
                                     CallingConv::ID CC,
                                     EVT VT) const override;
   unsigned getNumRegistersForCallingConv(LLVMContext &Context,
                                          CallingConv::ID CC,
                                          EVT VT) const override;

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

 #include "SIISelLowering.h"
 #include "AMDGPU.h"
 #include "AMDGPUInstrInfo.h"
 #include "AMDGPUTargetMachine.h"
 #include "SIMachineFunctionInfo.h"
 #include "SIRegisterInfo.h"
 #include "llvm/ADT/Statistic.h"
 #include "llvm/Analysis/LegacyDivergenceAnalysis.h"
+#include "llvm/Analysis/OptimizationRemarkEmitter.h"
 #include "llvm/BinaryFormat/ELF.h"
 #include "llvm/CodeGen/Analysis.h"
 #include "llvm/CodeGen/FunctionLoweringInfo.h"
 #include "llvm/CodeGen/GlobalISel/GISelKnownBits.h"
 #include "llvm/CodeGen/MachineLoopInfo.h"
 #include "llvm/IR/DiagnosticInfo.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/IntrinsicsAMDGPU.h"

 static bool fpModeMatchesGlobalFPAtomicMode(const AtomicRMWInst *RMW) {
   const fltSemantics &Flt = RMW->getType()->getScalarType()->getFltSemantics();
   auto DenormMode = RMW->getParent()->getParent()->getDenormalMode(Flt);
   if (&Flt == &APFloat::IEEEsingle())
     return DenormMode == DenormalMode::getPreserveSign();
   return DenormMode == DenormalMode::getIEEE();
 }
 
+static TargetLowering::AtomicExpansionKind
+atomicExpandReturn(OptimizationRemarkEmitter *ORE, AtomicRMWInst *RMW,
+                   TargetLowering::AtomicExpansionKind Kind, bool UnsafeFlag) {
+  ORE = new OptimizationRemarkEmitter(RMW->getFunction());
+  if (Kind == TargetLowering::AtomicExpansionKind::CmpXChg) {
+    ORE->emit([&]() {
+      OptimizationRemark Remark(DEBUG_TYPE, "Passed", RMW->getFunction());
+      Remark << "An FP atomic instruction was expanded into a CAS loop.";
+      return Remark;
+    });
+  } else if (Kind == TargetLowering::AtomicExpansionKind::None && UnsafeFlag) {
+    ORE->emit([&]() {
+      OptimizationRemark Remark(DEBUG_TYPE, "Passed", RMW->getFunction());
+      Remark << "An unsafe hardware instruction was generated.";
+      return Remark;
+    });
+  }
+  delete ORE;
+  return Kind;
+}
+
 TargetLowering::AtomicExpansionKind
 SITargetLowering::shouldExpandAtomicRMWInIR(AtomicRMWInst *RMW) const {
   switch (RMW->getOperation()) {
   case AtomicRMWInst::FAdd: {
     Type *Ty = RMW->getType();
 
     // We don't have a way to support 16-bit atomics now, so just leave them
     // as-is.
     if (Ty->isHalfTy())
       return AtomicExpansionKind::None;
 
     if (!Ty->isFloatTy() && (!Subtarget->hasGFX90AInsts() || !Ty->isDoubleTy()))
       return AtomicExpansionKind::CmpXChg;
 
     unsigned AS = RMW->getPointerAddressSpace();
-
+    bool UnsafeFPAtomicFlag = RMW->getFunction()
+                                  ->getFnAttribute("amdgpu-unsafe-fp-atomics")
+                                  .getValueAsBool();
     if ((AS == AMDGPUAS::GLOBAL_ADDRESS || AS == AMDGPUAS::FLAT_ADDRESS) &&
          Subtarget->hasAtomicFaddInsts()) {
       // The amdgpu-unsafe-fp-atomics attribute enables generation of unsafe
       // floating point atomic instructions. May generate more efficient code,
       // but may not respect rounding and denormal modes, and may give incorrect
       // results for certain memory destinations.
-      if (RMW->getFunction()
-              ->getFnAttribute("amdgpu-unsafe-fp-atomics")
-              .getValueAsString() != "true")
-        return AtomicExpansionKind::CmpXChg;
+      if (!UnsafeFPAtomicFlag)
+        return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+                                  UnsafeFPAtomicFlag);
+      atomicExpandReturn(ORE, RMW, AtomicExpansionKind::None, 
+        UnsafeFPAtomicFlag);
 
       if (Subtarget->hasGFX90AInsts()) {
         if (Ty->isFloatTy() && AS == AMDGPUAS::FLAT_ADDRESS)
-          return AtomicExpansionKind::CmpXChg;
+          return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+                                    UnsafeFPAtomicFlag);
 
         auto SSID = RMW->getSyncScopeID();
         if (SSID == SyncScope::System ||
             SSID == RMW->getContext().getOrInsertSyncScopeID("one-as"))
-          return AtomicExpansionKind::CmpXChg;
+          return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+                                    UnsafeFPAtomicFlag);
 
-        return AtomicExpansionKind::None;
+        return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::None,
+                                  UnsafeFPAtomicFlag);
       }
 
       if (AS == AMDGPUAS::FLAT_ADDRESS)
-        return AtomicExpansionKind::CmpXChg;
+        return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+                                  UnsafeFPAtomicFlag);
 
-      return RMW->use_empty() ? AtomicExpansionKind::None
+      auto Kind = RMW->use_empty() ? AtomicExpansionKind::None
                                    : AtomicExpansionKind::CmpXChg;
+      return atomicExpandReturn(ORE, RMW, Kind, UnsafeFPAtomicFlag);
     }
 
     // DS FP atomics do repect the denormal mode, but the rounding mode is fixed
     // to round-to-nearest-even.
     // The only exception is DS_ADD_F64 which never flushes regardless of mode.
     if (AS == AMDGPUAS::LOCAL_ADDRESS && Subtarget->hasLDSFPAtomics()) {
       if (!Ty->isDoubleTy())
-        return AtomicExpansionKind::None;
+        return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::None,
+                                  UnsafeFPAtomicFlag);
 
-      return (fpModeMatchesGlobalFPAtomicMode(RMW) ||
-              RMW->getFunction()
-                      ->getFnAttribute("amdgpu-unsafe-fp-atomics")
-                      .getValueAsString() == "true")
+      auto Kind = (fpModeMatchesGlobalFPAtomicMode(RMW) || UnsafeFPAtomicFlag)
                       ? AtomicExpansionKind::None
                       : AtomicExpansionKind::CmpXChg;
+      return atomicExpandReturn(ORE, RMW, Kind, UnsafeFPAtomicFlag);
     }
 
-    return AtomicExpansionKind::CmpXChg;
+    return atomicExpandReturn(ORE, RMW, AtomicExpansionKind::CmpXChg,
+                              UnsafeFPAtomicFlag);
   }
   default:
     break;
   }
 
   return AMDGPUTargetLowering::shouldExpandAtomicRMWInIR(RMW);
 }