Skip to content

Commit dd352ab

Browse files
arsenmarsenm
committed
clang/AMDGPU: Emit atomicrmw for global/flat fadd v2bf16 builtins
1 parent 0a22655 commit dd352ab

File tree

3 files changed

+30
-24
lines changed

3 files changed

+30
-24
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 8 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -18955,22 +18955,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1895518955
CGM.getIntrinsic(IID, {ArgTy, Addr->getType(), Val->getType()});
1895618956
return Builder.CreateCall(F, {Addr, Val});
1895718957
}
18958-
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
18959-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: {
18960-
Intrinsic::ID IID;
18961-
switch (BuiltinID) {
18962-
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
18963-
IID = Intrinsic::amdgcn_global_atomic_fadd_v2bf16;
18964-
break;
18965-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
18966-
IID = Intrinsic::amdgcn_flat_atomic_fadd_v2bf16;
18967-
break;
18968-
}
18969-
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
18970-
llvm::Value *Val = EmitScalarExpr(E->getArg(1));
18971-
llvm::Function *F = CGM.getIntrinsic(IID, {Addr->getType()});
18972-
return Builder.CreateCall(F, {Addr, Val});
18973-
}
1897418958
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
1897518959
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
1897618960
case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
@@ -19352,7 +19336,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1935219336
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
1935319337
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
1935419338
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
19355-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64: {
19339+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
19340+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
19341+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: {
1935619342
llvm::AtomicRMWInst::BinOp BinOp;
1935719343
switch (BuiltinID) {
1935819344
case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
@@ -19374,6 +19360,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1937419360
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
1937519361
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
1937619362
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
19363+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
19364+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
1937719365
BinOp = llvm::AtomicRMWInst::FAdd;
1937819366
break;
1937919367
case AMDGPU::BI__builtin_amdgcn_ds_fminf:
@@ -19418,7 +19406,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1941819406
AO = AtomicOrdering::Monotonic;
1941919407

1942019408
// The v2bf16 builtin uses i16 instead of a natural bfloat type.
19421-
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16) {
19409+
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16 ||
19410+
BuiltinID == AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16 ||
19411+
BuiltinID == AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16) {
1942219412
llvm::Type *V2BF16Ty = FixedVectorType::get(
1942319413
llvm::Type::getBFloatTy(Builder.getContext()), 2);
1942419414
Val = Builder.CreateBitCast(Val, V2BF16Ty);

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

Lines changed: 14 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@ typedef short __attribute__((ext_vector_type(2))) short2;
1111

1212
// CHECK-LABEL: test_local_add_2bf16
1313
// CHECK: [[BC0:%.+]] = bitcast <2 x i16> {{.+}} to <2 x bfloat>
14-
// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x bfloat> [[BC0]] syncscope("agent") monotonic, align 4
14+
// CHECK-NEXT: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(3) %{{.+}}, <2 x bfloat> [[BC0]] syncscope("agent") monotonic, align 4
1515
// CHECK-NEXT: bitcast <2 x bfloat> [[RMW]] to <2 x i16>
1616

1717
// GFX12-LABEL: test_local_add_2bf16
@@ -57,7 +57,10 @@ half2 test_flat_add_2f16(__generic half2 *addr, half2 x) {
5757
}
5858

5959
// CHECK-LABEL: test_flat_add_2bf16
60-
// CHECK: call <2 x i16> @llvm.amdgcn.flat.atomic.fadd.v2bf16.p0(ptr %{{.*}}, <2 x i16> %{{.*}})
60+
// CHECK: [[BC:%.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat>
61+
// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr %{{.+}}, <2 x bfloat> [[BC]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
62+
// CHECK: bitcast <2 x bfloat> [[RMW]] to <2 x i16>
63+
6164
// GFX12-LABEL: test_flat_add_2bf16
6265
// GFX12: flat_atomic_pk_add_bf16
6366
short2 test_flat_add_2bf16(__generic short2 *addr, short2 x) {
@@ -84,7 +87,11 @@ void test_global_add_half2_noret(__global half2 *addr, half2 x) {
8487
}
8588

8689
// CHECK-LABEL: test_global_add_2bf16
87-
// CHECK: call <2 x i16> @llvm.amdgcn.global.atomic.fadd.v2bf16.p1(ptr addrspace(1) %{{.*}}, <2 x i16> %{{.*}})
90+
// CHECK: [[BC:%.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat>
91+
// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x bfloat> [[BC]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
92+
// CHECK: bitcast <2 x bfloat> [[RMW]] to <2 x i16>
93+
94+
8895
// GFX12-LABEL: test_global_add_2bf16
8996
// GFX12: global_atomic_pk_add_bf16 v2, v[0:1], v2, off th:TH_ATOMIC_RETURN
9097
void test_global_add_2bf16(__global short2 *addr, short2 x) {
@@ -93,7 +100,10 @@ void test_global_add_2bf16(__global short2 *addr, short2 x) {
93100
}
94101

95102
// CHECK-LABEL: test_global_add_2bf16_noret
96-
// CHECK: call <2 x i16> @llvm.amdgcn.global.atomic.fadd.v2bf16.p1(ptr addrspace(1) %{{.*}}, <2 x i16> %{{.*}})
103+
// CHECK: [[BC:%.+]] = bitcast <2 x i16> %{{.+}} to <2 x bfloat>
104+
// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x bfloat> [[BC]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
105+
// CHECK: bitcast <2 x bfloat> [[RMW]] to <2 x i16>
106+
97107
// GFX12-LABEL: test_global_add_2bf16_noret
98108
// GFX12: global_atomic_pk_add_bf16 v[0:1], v2, off
99109
void test_global_add_2bf16_noret(__global short2 *addr, short2 x) {

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

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -28,15 +28,21 @@ half2 test_flat_add_2f16(__generic half2 *addr, half2 x) {
2828
}
2929

3030
// CHECK-LABEL: test_flat_add_2bf16
31-
// CHECK: call <2 x i16> @llvm.amdgcn.flat.atomic.fadd.v2bf16.p0(ptr %{{.*}}, <2 x i16> %{{.*}})
31+
// CHECK: [[BC0:%.+]] = bitcast <2 x i16> {{.+}} to <2 x bfloat>
32+
// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr %{{.+}}, <2 x bfloat> [[BC0]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
33+
// CHECK-NEXT: bitcast <2 x bfloat> [[RMW]] to <2 x i16>
34+
3235
// GFX940-LABEL: test_flat_add_2bf16
3336
// GFX940: flat_atomic_pk_add_bf16
3437
short2 test_flat_add_2bf16(__generic short2 *addr, short2 x) {
3538
return __builtin_amdgcn_flat_atomic_fadd_v2bf16(addr, x);
3639
}
3740

3841
// CHECK-LABEL: test_global_add_2bf16
39-
// CHECK: call <2 x i16> @llvm.amdgcn.global.atomic.fadd.v2bf16.p1(ptr addrspace(1) %{{.*}}, <2 x i16> %{{.*}})
42+
// CHECK: [[BC0:%.+]] = bitcast <2 x i16> {{.+}} to <2 x bfloat>
43+
// CHECK: [[RMW:%.+]] = atomicrmw fadd ptr addrspace(1) %{{.+}}, <2 x bfloat> [[BC0]] syncscope("agent") monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
44+
// CHECK-NEXT: bitcast <2 x bfloat> [[RMW]] to <2 x i16>
45+
4046
// GFX940-LABEL: test_global_add_2bf16
4147
// GFX940: global_atomic_pk_add_bf16
4248
short2 test_global_add_2bf16(__global short2 *addr, short2 x) {

0 commit comments

Comments
 (0)