Skip to content

Commit 6e2f756

Browse files
arsenmarsenm
committed
clang/AMDGPU: Emit atomicrmw for flat/global atomic min/max f64 builtins
1 parent 16b9d36 commit 6e2f756

File tree

2 files changed

+21
-33
lines changed

2 files changed

+21
-33
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 9 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -18633,32 +18633,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1863318633
Function *F = CGM.getIntrinsic(Intrin, { Src0->getType() });
1863418634
return Builder.CreateCall(F, { Src0, Builder.getFalse() });
1863518635
}
18636-
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
18637-
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
18638-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
18639-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
18640-
Intrinsic::ID IID;
18641-
llvm::Type *ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
18642-
switch (BuiltinID) {
18643-
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
18644-
IID = Intrinsic::amdgcn_global_atomic_fmin;
18645-
break;
18646-
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
18647-
IID = Intrinsic::amdgcn_global_atomic_fmax;
18648-
break;
18649-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
18650-
IID = Intrinsic::amdgcn_flat_atomic_fmin;
18651-
break;
18652-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:
18653-
IID = Intrinsic::amdgcn_flat_atomic_fmax;
18654-
break;
18655-
}
18656-
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
18657-
llvm::Value *Val = EmitScalarExpr(E->getArg(1));
18658-
llvm::Function *F =
18659-
CGM.getIntrinsic(IID, {ArgTy, Addr->getType(), Val->getType()});
18660-
return Builder.CreateCall(F, {Addr, Val});
18661-
}
1866218636
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
1866318637
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
1866418638
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
@@ -19034,7 +19008,11 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1903419008
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
1903519009
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
1903619010
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
19037-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: {
19011+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
19012+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
19013+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
19014+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
19015+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
1903819016
llvm::AtomicRMWInst::BinOp BinOp;
1903919017
switch (BuiltinID) {
1904019018
case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
@@ -19061,8 +19039,12 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1906119039
BinOp = llvm::AtomicRMWInst::FAdd;
1906219040
break;
1906319041
case AMDGPU::BI__builtin_amdgcn_ds_fminf:
19042+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
19043+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
1906419044
BinOp = llvm::AtomicRMWInst::FMin;
1906519045
break;
19046+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
19047+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:
1906619048
case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
1906719049
BinOp = llvm::AtomicRMWInst::FMax;
1906819050
break;

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

Lines changed: 12 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,8 @@ void test_global_add_half2(__global half2 *addr, half2 x) {
2727
}
2828

2929
// CHECK-LABEL: test_global_global_min_f64
30-
// CHECK: call double @llvm.amdgcn.global.atomic.fmin.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}})
30+
// CHECK: = atomicrmw fmin ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
31+
3132
// GFX90A-LABEL: test_global_global_min_f64$local
3233
// GFX90A: global_atomic_min_f64
3334
void test_global_global_min_f64(__global double *addr, double x){
@@ -36,7 +37,8 @@ void test_global_global_min_f64(__global double *addr, double x){
3637
}
3738

3839
// CHECK-LABEL: test_global_max_f64
39-
// CHECK: call double @llvm.amdgcn.global.atomic.fmax.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}})
40+
// CHECK: = atomicrmw fmax ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
41+
4042
// GFX90A-LABEL: test_global_max_f64$local
4143
// GFX90A: global_atomic_max_f64
4244
void test_global_max_f64(__global double *addr, double x){
@@ -65,7 +67,8 @@ void test_flat_global_add_f64(__global double *addr, double x){
6567
}
6668

6769
// CHECK-LABEL: test_flat_min_flat_f64
68-
// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p0.f64(ptr %{{.*}}, double %{{.*}})
70+
// CHECK: = atomicrmw fmin ptr {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
71+
6972
// GFX90A-LABEL: test_flat_min_flat_f64$local
7073
// GFX90A: flat_atomic_min_f64
7174
void test_flat_min_flat_f64(__generic double *addr, double x){
@@ -74,7 +77,8 @@ void test_flat_min_flat_f64(__generic double *addr, double x){
7477
}
7578

7679
// CHECK-LABEL: test_flat_global_min_f64
77-
// CHECK: call double @llvm.amdgcn.flat.atomic.fmin.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}})
80+
// CHECK: = atomicrmw fmin ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
81+
7882
// GFX90A: test_flat_global_min_f64$local
7983
// GFX90A: global_atomic_min_f64
8084
void test_flat_global_min_f64(__global double *addr, double x){
@@ -83,7 +87,8 @@ void test_flat_global_min_f64(__global double *addr, double x){
8387
}
8488

8589
// CHECK-LABEL: test_flat_max_flat_f64
86-
// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p0.f64(ptr %{{.*}}, double %{{.*}})
90+
// CHECK: = atomicrmw fmax ptr {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
91+
8792
// GFX90A-LABEL: test_flat_max_flat_f64$local
8893
// GFX90A: flat_atomic_max_f64
8994
void test_flat_max_flat_f64(__generic double *addr, double x){
@@ -92,7 +97,8 @@ void test_flat_max_flat_f64(__generic double *addr, double x){
9297
}
9398

9499
// CHECK-LABEL: test_flat_global_max_f64
95-
// CHECK: call double @llvm.amdgcn.flat.atomic.fmax.f64.p1.f64(ptr addrspace(1) %{{.*}}, double %{{.*}})
100+
// CHECK: = atomicrmw fmax ptr addrspace(1) {{.+}}, double %{{.+}} syncscope("agent") monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
101+
96102
// GFX90A-LABEL: test_flat_global_max_f64$local
97103
// GFX90A: global_atomic_max_f64
98104
void test_flat_global_max_f64(__global double *addr, double x){

0 commit comments

Comments
 (0)