[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

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

clang/lib/CodeGen/CGBuiltin.cpp

16247

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

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

112 ↗

(On Diff #362804)

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

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

16247

You mean push it down after line 16119?

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

112 ↗

(On Diff #362804)

Sorry, what do you mean by rtn version?

[arsenm]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201

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

16247

Yes

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

112 ↗

(On Diff #362804)

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

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

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

Flat address space is 0.

210

This is available since gfx8. Attribute gfx8-insts.

clang/lib/CodeGen/CGBuiltin.cpp

16208

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

112 ↗

(On Diff #362804)

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

[gandhi21299]

clang/lib/CodeGen/CGBuiltin.cpp

16208

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

16208

Every instantiation of GCCBuiltin in the IntrinsicsAMDGPU.td.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16208

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

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16208

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

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16208

These intrinsics are all mangled based on the FP type

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16208

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

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16266

Should we map flags since we already have them?

[gandhi21299]

clang/lib/CodeGen/CGBuiltin.cpp

16266

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

16266

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

16266

Following the discussion, what change is required here?

[rampitec]

clang/lib/CodeGen/CGBuiltin.cpp

16266

Keep zeroes, drop immediate argument of the builtins.

[gandhi21299]

updating test

[gandhi21299]

clang/include/clang/Basic/BuiltinsAMDGPU.def

201

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

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

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

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

14

Use _f64 or _double in the test name.

32

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

68

'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/include/clang/Basic/BuiltinsAMDGPU.def

 TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi", "nc", "gfx8-insts")
 
 //===----------------------------------------------------------------------===//
 // GFX9+ only builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
 
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2hV2h*1V2h", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1d", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmax_f64, "dd*1d", "t", "gfx90a-insts")
+
+TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f64, "dd*0d", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmin_f64, "dd*0d", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d", "t", "gfx90a-insts")
+
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t", "gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
+
 //===----------------------------------------------------------------------===//
 // Deep learning builtins.
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb", "nc", "dot7-insts")
 TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb", "nc", "dot2-insts")
 TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb", "nc", "dot2-insts")
 TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb", "nc", "dot1-insts")

clang/lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: {
     llvm::Value *Src3 = EmitScalarExpr(E->getArg(3));
     llvm::Value *Src4 = EmitScalarExpr(E->getArg(4));
     llvm::Function *F = CGM.getIntrinsic(Intrin, { Src1->getType() });
     llvm::FunctionType *FTy = F->getFunctionType();
     llvm::Type *PTy = FTy->getParamType(0);
     Src0 = Builder.CreatePointerBitCastOrAddrSpaceCast(Src0, PTy);
     return Builder.CreateCall(F, { Src0, Src1, Src2, Src3, Src4 });
   }
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
+  case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
+  case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
+  case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
+  case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
+    Intrinsic::ID IID;
+    llvm::Type *ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
+    switch (BuiltinID) {
+    case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
+      ArgTy = llvm::Type::getFloatTy(getLLVMContext());
+      IID = Intrinsic::amdgcn_global_atomic_fadd;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
+      ArgTy = llvm::FixedVectorType::get(
+          llvm::Type::getHalfTy(getLLVMContext()), 2);
+      IID = Intrinsic::amdgcn_global_atomic_fadd;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
+      IID = Intrinsic::amdgcn_global_atomic_fadd;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
+      IID = Intrinsic::amdgcn_global_atomic_fmin;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
+      IID = Intrinsic::amdgcn_global_atomic_fmax;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
+      IID = Intrinsic::amdgcn_flat_atomic_fadd;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
+      IID = Intrinsic::amdgcn_flat_atomic_fmin;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:
+      IID = Intrinsic::amdgcn_flat_atomic_fmax;
+      break;
+    }
+    llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
+    llvm::Value *Val = EmitScalarExpr(E->getArg(1));
+    llvm::Function *F =
+        CGM.getIntrinsic(IID, {ArgTy, Addr->getType(), Val->getType()});
+    return Builder.CreateCall(F, {Addr, Val});
+  }
+  case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
+  case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32: {
+    Intrinsic::ID IID;
+    llvm::Type *ArgTy;
+    switch (BuiltinID) {
+    case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
+      ArgTy = llvm::Type::getFloatTy(getLLVMContext());
+      IID = Intrinsic::amdgcn_ds_fadd;
+      break;
+    case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
+      ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
+      IID = Intrinsic::amdgcn_ds_fadd;
+      break;
+    }
+    llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
+    llvm::Value *Val = EmitScalarExpr(E->getArg(1));
+    llvm::Constant *ZeroI32 = llvm::ConstantInt::getIntegerValue(
+        llvm::Type::getInt32Ty(getLLVMContext()), APInt(32, 0, true));
+    llvm::Constant *ZeroI1 = llvm::ConstantInt::getIntegerValue(
+        llvm::Type::getInt1Ty(getLLVMContext()), APInt(1, 0));
+    llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy});
+    return Builder.CreateCall(F, {Addr, Val, ZeroI32, ZeroI32, ZeroI1});
+  }
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, NormalRead, "exec"));
     CI->setConvergent();
     return CI;
   }
   case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
   case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {

clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx7-err.cl

+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx700 \
+// RUN:   %s -verify -S -o -
+
+kernel void test_fadd_local(__local float *ptr, float val){
+    float *res;
+    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val); // expected-error{{'__builtin_amdgcn_ds_atomic_fadd_f32' needs target feature gfx8-insts}}
+}

clang/test/CodeGenOpenCL/builtins-amdgcn-fp-atomics-gfx908-err.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx908 \
+// RUN:   -verify -S -o - %s
+
+typedef half __attribute__((ext_vector_type(2))) half2;
+
+void test_global_add_2f16(__global half2 *addrh2, half2 xh2,
+                          __global float *addrf, float xf,
+                          __global double *addr, double x) {
+  half2 *half_rtn;
+  float *fp_rtn;
+  double *rtn;
+  *half_rtn = __builtin_amdgcn_global_atomic_fadd_v2f16(addrh2, xh2); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2f16' needs target feature gfx90a-insts}}
+  *fp_rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_f32' needs target feature gfx90a-insts}}
+  *rtn = __builtin_amdgcn_global_atomic_fadd_f64(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_f64' needs target feature gfx90a-insts}}
+  *rtn = __builtin_amdgcn_global_atomic_fmax_f64(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fmax_f64' needs target feature gfx90a-insts}}
+  *rtn = __builtin_amdgcn_global_atomic_fmin_f64(addr, x); // expected-error{{'__builtin_amdgcn_global_atomic_fmin_f64' needs target feature gfx90a-insts}}
+  *rtn = __builtin_amdgcn_flat_atomic_fadd_f64(addr, x); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f64' needs target feature gfx90a-insts}}
+  *rtn = __builtin_amdgcn_flat_atomic_fmin_f64(addr, x); // expected-error{{'__builtin_amdgcn_flat_atomic_fmin_f64' needs target feature gfx90a-insts}}
+  *rtn = __builtin_amdgcn_flat_atomic_fmax_f64(addr, x); // expected-error{{'__builtin_amdgcn_flat_atomic_fmax_f64' needs target feature gfx90a-insts}}
+}

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

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \
+// RUN:   -S -o - %s
+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1030 \
+// RUN:   -S -o - %s | FileCheck -check-prefix=GFX1030 %s
+
+// CHECK-LABEL: test_ds_addf_local
+// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}},
+// GFX1030-LABEL:  test_ds_addf_local$local
+// GFX1030:  ds_add_rtn_f32
+void test_ds_addf_local(__local float *addr, float x){
+  float *rtn;
+  *rtn = __builtin_amdgcn_ds_atomic_fadd_f32(addr, x);
+}

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

+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
+// RUN:   %s -S -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx810 \
+// RUN:   -S -o - %s | FileCheck -check-prefix=GFX8 %s
+
+// CHECK-LABEL: test_fadd_local
+// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}}, i32 0, i32 0, i1 false)
+// GFX8-LABEL: test_fadd_local$local:
+// GFX8: ds_add_rtn_f32 v2, v0, v1
+// GFX8: s_endpgm
+kernel void test_fadd_local(__local float *ptr, float val){
+    float *res;
+    *res = __builtin_amdgcn_ds_atomic_fadd_f32(ptr, val);
+}

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

+// RUN: %clang_cc1 -O0 -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 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN:   -S -o - %s | FileCheck -check-prefix=GFX90A %s
+
+typedef half __attribute__((ext_vector_type(2))) half2;
+
+// CHECK-LABEL: test_global_add_f64
+// CHECK: call double @llvm.amdgcn.global.atomic.fadd.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_global_add_f64$local:
+// GFX90A:  global_atomic_add_f64
+void test_global_add_f64(__global double *addr, double x) {
+  double *rtn;
+  *rtn = __builtin_amdgcn_global_atomic_fadd_f64(addr, x);
+}
+
+// CHECK-LABEL: test_global_add_half2
+// CHECK: call <2 x half> @llvm.amdgcn.global.atomic.fadd.v2f16.p1v2f16.v2f16(<2 x half> addrspace(1)* %{{.*}}, <2 x half> %{{.*}})
+// GFX90A-LABEL:  test_global_add_half2
+// GFX90A:  global_atomic_pk_add_f16 v2, v[0:1], v2, off glc
+void test_global_add_half2(__global half2 *addr, half2 x) {
+  half2 *rtn;
+  *rtn = __builtin_amdgcn_global_atomic_fadd_v2f16(addr, x);
+}
+
+// CHECK-LABEL: test_global_global_min_f64
+// CHECK: call double @llvm.amdgcn.global.atomic.fmin.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_global_global_min_f64$local
+// GFX90A:  global_atomic_min_f64
+void test_global_global_min_f64(__global double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_global_atomic_fmin_f64(addr, x);
+}
+
+// CHECK-LABEL: test_global_max_f64
+// CHECK: call double @llvm.amdgcn.global.atomic.fmax.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_global_max_f64$local
+// GFX90A:  global_atomic_max_f64
+void test_global_max_f64(__global double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_global_atomic_fmax_f64(addr, x);
+}
+
+// CHECK-LABEL: test_flat_add_local_f64
+// CHECK: call double @llvm.amdgcn.flat.atomic.fadd.f64.p3f64.f64(double addrspace(3)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_add_local_f64$local
+// GFX90A:  ds_add_rtn_f64
+void test_flat_add_local_f64(__local double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_flat_atomic_fadd_f64(addr, x);
+}
+
+// CHECK-LABEL: test_flat_global_add_f64
+// CHECK: call double @llvm.amdgcn.flat.atomic.fadd.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_global_add_f64$local
+// GFX90A:  global_atomic_add_f64
+void test_flat_global_add_f64(__global double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_flat_atomic_fadd_f64(addr, x);
+}
+
+// CHECK-LABEL: test_flat_min_flat_f64
+// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p0f64.f64(double* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_min_flat_f64$local
+// GFX90A:  flat_atomic_min_f64
+void test_flat_min_flat_f64(__generic double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_flat_atomic_fmin_f64(addr, x);
+}
+
+// CHECK-LABEL: test_flat_global_min_f64
+// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A:  test_flat_global_min_f64$local
+// GFX90A:  global_atomic_min_f64
+void test_flat_global_min_f64(__global double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_flat_atomic_fmin_f64(addr, x);
+}
+
+// CHECK-LABEL: test_flat_max_flat_f64
+// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p0f64.f64(double* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_max_flat_f64$local
+// GFX90A:  flat_atomic_max_f64
+void test_flat_max_flat_f64(__generic double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_flat_atomic_fmax_f64(addr, x);
+}
+
+// CHECK-LABEL: test_flat_global_max_f64
+// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p1f64.f64(double addrspace(1)* %{{.*}}, double %{{.*}})
+// GFX90A-LABEL:  test_flat_global_max_f64$local
+// GFX90A:  global_atomic_max_f64
+void test_flat_global_max_f64(__global double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_flat_atomic_fmax_f64(addr, x);
+}
+
+// CHECK-LABEL: test_ds_add_local_f64
+// CHECK: call double @llvm.amdgcn.ds.fadd.f64(double addrspace(3)* %{{.*}}, double %{{.*}},
+// GFX90A:  test_ds_add_local_f64$local
+// GFX90A:  ds_add_rtn_f64
+void test_ds_add_local_f64(__local double *addr, double x){
+  double *rtn;
+  *rtn = __builtin_amdgcn_ds_atomic_fadd_f64(addr, x);
+}
+
+// CHECK-LABEL: test_ds_addf_local_f32
+// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %{{.*}}, float %{{.*}},
+// GFX90A-LABEL:  test_ds_addf_local_f32$local
+// GFX90A:  ds_add_rtn_f32
+void test_ds_addf_local_f32(__local float *addr, float x){
+  float *rtn;
+  *rtn = __builtin_amdgcn_ds_atomic_fadd_f32(addr, x);
+}