Skip to content

Commit 932f628

Browse files
rampitecrampitec
committed
[AMDGPU] new gfx940 fp atomics
Differential Revision: https://reviews.llvm.org/D121028
1 parent 89d5c31 commit 932f628

File tree

16 files changed

+790
-4
lines changed

16 files changed

+790
-4
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -209,6 +209,12 @@ TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d", "t", "gfx90a-inst
209209
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d", "t", "gfx90a-insts")
210210
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f", "t", "gfx8-insts")
211211

212+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f", "t", "gfx940-insts")
213+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h", "t", "gfx940-insts")
214+
TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s", "t", "gfx940-insts")
215+
TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "gfx940-insts")
216+
TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "gfx940-insts")
217+
212218
//===----------------------------------------------------------------------===//
213219
// Deep learning builtins.
214220
//===----------------------------------------------------------------------===//

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 28 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -16513,7 +16513,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1651316513
case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
1651416514
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
1651516515
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
16516-
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
16516+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:
16517+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
16518+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: {
1651716519
Intrinsic::ID IID;
1651816520
llvm::Type *ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
1651916521
switch (BuiltinID) {
@@ -16544,13 +16546,38 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
1654416546
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:
1654516547
IID = Intrinsic::amdgcn_flat_atomic_fmax;
1654616548
break;
16549+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
16550+
ArgTy = llvm::Type::getFloatTy(getLLVMContext());
16551+
IID = Intrinsic::amdgcn_flat_atomic_fadd;
16552+
break;
16553+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
16554+
ArgTy = llvm::FixedVectorType::get(
16555+
llvm::Type::getHalfTy(getLLVMContext()), 2);
16556+
IID = Intrinsic::amdgcn_flat_atomic_fadd;
16557+
break;
1654716558
}
1654816559
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
1654916560
llvm::Value *Val = EmitScalarExpr(E->getArg(1));
1655016561
llvm::Function *F =
1655116562
CGM.getIntrinsic(IID, {ArgTy, Addr->getType(), Val->getType()});
1655216563
return Builder.CreateCall(F, {Addr, Val});
1655316564
}
16565+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
16566+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: {
16567+
Intrinsic::ID IID;
16568+
switch (BuiltinID) {
16569+
case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
16570+
IID = Intrinsic::amdgcn_global_atomic_fadd_v2bf16;
16571+
break;
16572+
case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
16573+
IID = Intrinsic::amdgcn_flat_atomic_fadd_v2bf16;
16574+
break;
16575+
}
16576+
llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
16577+
llvm::Value *Val = EmitScalarExpr(E->getArg(1));
16578+
llvm::Function *F = CGM.getIntrinsic(IID, {Addr->getType()});
16579+
return Builder.CreateCall(F, {Addr, Val});
16580+
}
1655416581
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
1655516582
case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32: {
1655616583
Intrinsic::ID IID;

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

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
2+
// RUN: -verify -S -o - %s
3+
4+
// REQUIRES: amdgpu-registered-target
5+
6+
typedef half __attribute__((ext_vector_type(2))) half2;
7+
typedef short __attribute__((ext_vector_type(2))) short2;
8+
9+
void test_atomic_fadd(__global half2 *addrh2, half2 xh2,
10+
__global short2 *addrs2, __local short2 *addrs2l, short2 xs2,
11+
__global float *addrf, float xf) {
12+
__builtin_amdgcn_flat_atomic_fadd_f32(addrf, xf); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_f32' needs target feature gfx940-insts}}
13+
__builtin_amdgcn_flat_atomic_fadd_v2f16(addrh2, xh2); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2f16' needs target feature gfx940-insts}}
14+
__builtin_amdgcn_flat_atomic_fadd_v2bf16(addrs2, xs2); // expected-error{{'__builtin_amdgcn_flat_atomic_fadd_v2bf16' needs target feature gfx940-insts}}
15+
__builtin_amdgcn_global_atomic_fadd_v2bf16(addrs2, xs2); // expected-error{{'__builtin_amdgcn_global_atomic_fadd_v2bf16' needs target feature gfx940-insts}}
16+
__builtin_amdgcn_ds_atomic_fadd_v2bf16(addrs2l, xs2); // expected-error{{'__builtin_amdgcn_ds_atomic_fadd_v2bf16' needs target feature gfx940-insts}}
17+
}

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

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,50 @@
1+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 \
2+
// RUN: %s -S -emit-llvm -o - | FileCheck %s
3+
4+
// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx940 \
5+
// RUN: -S -o - %s | FileCheck -check-prefix=GFX940 %s
6+
7+
// REQUIRES: amdgpu-registered-target
8+
9+
typedef half __attribute__((ext_vector_type(2))) half2;
10+
typedef short __attribute__((ext_vector_type(2))) short2;
11+
12+
// CHECK-LABEL: test_flat_add_f32
13+
// CHECK: call float @llvm.amdgcn.flat.atomic.fadd.f32.p0f32.f32(float* %{{.*}}, float %{{.*}})
14+
// GFX940-LABEL: test_flat_add_f32
15+
// GFX940: flat_atomic_add_f32
16+
half2 test_flat_add_f32(__generic float *addr, float x) {
17+
return __builtin_amdgcn_flat_atomic_fadd_f32(addr, x);
18+
}
19+
20+
// CHECK-LABEL: test_flat_add_2f16
21+
// CHECK: call <2 x half> @llvm.amdgcn.flat.atomic.fadd.v2f16.p0v2f16.v2f16(<2 x half>* %{{.*}}, <2 x half> %{{.*}})
22+
// GFX940-LABEL: test_flat_add_2f16
23+
// GFX940: flat_atomic_pk_add_f16
24+
half2 test_flat_add_2f16(__generic half2 *addr, half2 x) {
25+
return __builtin_amdgcn_flat_atomic_fadd_v2f16(addr, x);
26+
}
27+
28+
// CHECK-LABEL: test_flat_add_2bf16
29+
// CHECK: call <2 x i16> @llvm.amdgcn.flat.atomic.fadd.v2bf16.p0v2i16(<2 x i16>* %{{.*}}, <2 x i16> %{{.*}})
30+
// GFX940-LABEL: test_flat_add_2bf16
31+
// GFX940: flat_atomic_pk_add_bf16
32+
short2 test_flat_add_2bf16(__generic short2 *addr, short2 x) {
33+
return __builtin_amdgcn_flat_atomic_fadd_v2bf16(addr, x);
34+
}
35+
36+
// CHECK-LABEL: test_global_add_2bf16
37+
// CHECK: call <2 x i16> @llvm.amdgcn.global.atomic.fadd.v2bf16.p1v2i16(<2 x i16> addrspace(1)* %{{.*}}, <2 x i16> %{{.*}})
38+
// GFX940-LABEL: test_global_add_2bf16
39+
// GFX940: global_atomic_pk_add_bf16
40+
short2 test_global_add_2bf16(__global short2 *addr, short2 x) {
41+
return __builtin_amdgcn_global_atomic_fadd_v2bf16(addr, x);
42+
}
43+
44+
// CHECK-LABEL: test_local_add_2bf16
45+
// CHECK: call <2 x i16> @llvm.amdgcn.ds.fadd.v2bf16(<2 x i16> addrspace(3)* %{{.*}}, <2 x i16> %
46+
// GFX940-LABEL: test_local_add_2bf16
47+
// GFX940: ds_pk_add_rtn_bf16
48+
short2 test_local_add_2bf16(__local short2 *addr, short2 x) {
49+
return __builtin_amdgcn_ds_atomic_fadd_v2bf16(addr, x);
50+
}

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1978,6 +1978,19 @@ def int_amdgcn_mfma_f32_16x16x16bf16_1k : AMDGPUMfmaIntrinsic<llvm_v4f32_ty, ll
19781978
def int_amdgcn_mfma_f64_16x16x4f64 : AMDGPUMfmaIntrinsic<llvm_v4f64_ty, llvm_double_ty>;
19791979
def int_amdgcn_mfma_f64_4x4x4f64 : AMDGPUMfmaIntrinsic<llvm_double_ty, llvm_double_ty>;
19801980

1981+
//===----------------------------------------------------------------------===//
1982+
// gfx940 intrinsics
1983+
// ===----------------------------------------------------------------------===//
1984+
1985+
// bf16 atomics use v2i16 argument since there is no bf16 data type in the llvm.
1986+
def int_amdgcn_global_atomic_fadd_v2bf16 : AMDGPUGlobalAtomicRtn<llvm_v2i16_ty>;
1987+
def int_amdgcn_flat_atomic_fadd_v2bf16 : AMDGPUGlobalAtomicRtn<llvm_v2i16_ty>;
1988+
def int_amdgcn_ds_fadd_v2bf16 : Intrinsic<
1989+
[llvm_v2i16_ty],
1990+
[LLVMQualPointerType<llvm_v2i16_ty, 3>, llvm_v2i16_ty],
1991+
[IntrArgMemOnly, IntrWillReturn, NoCapture<ArgIndex<0>>]>,
1992+
GCCBuiltin<"__builtin_amdgcn_ds_atomic_fadd_v2bf16">;
1993+
19811994
//===----------------------------------------------------------------------===//
19821995
// Special Intrinsics for backend internal use only. No frontend
19831996
// should emit calls to these.

‎llvm/lib/Target/AMDGPU/AMDGPU.td‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1558,6 +1558,13 @@ def HasFmaLegacy32 : Predicate<"Subtarget->hasGFX10_3Insts()">,
15581558
def HasAtomicFaddInsts : Predicate<"Subtarget->hasAtomicFaddInsts()">,
15591559
AssemblerPredicate<(all_of FeatureAtomicFaddInsts)>;
15601560

1561+
// Differentiate between two functionally equivalent, but incompatible
1562+
// encoding-wise FP atomics between gfx90* and gfx940
1563+
def HasAtomicFaddInstsGFX90X : Predicate<"Subtarget->hasAtomicFaddInsts()">,
1564+
AssemblerPredicate<(all_of FeatureAtomicFaddInsts, (not FeatureGFX940Insts))>;
1565+
def HasAtomicFaddInstsGFX940 : Predicate<"Subtarget->hasAtomicFaddInsts()">,
1566+
AssemblerPredicate<(all_of FeatureAtomicFaddInsts, FeatureGFX940Insts)>;
1567+
15611568
def HasDsSrc2Insts : Predicate<"!Subtarget->hasDsSrc2Insts()">,
15621569
AssemblerPredicate<(all_of FeatureDsSrc2Insts)>;
15631570

‎llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1297,6 +1297,8 @@ AMDGPULegalizerInfo::AMDGPULegalizerInfo(const GCNSubtarget &ST_,
12971297
Atomic.legalFor({{S32, LocalPtr}, {S32, RegionPtr}});
12981298
if (ST.hasGFX90AInsts())
12991299
Atomic.legalFor({{S64, LocalPtr}});
1300+
if (ST.hasGFX940Insts())
1301+
Atomic.legalFor({{V2S16, LocalPtr}});
13001302
}
13011303
if (ST.hasAtomicFaddInsts())
13021304
Atomic.legalFor({{S32, GlobalPtr}});

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4339,6 +4339,8 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
43394339
case Intrinsic::amdgcn_flat_atomic_fadd:
43404340
case Intrinsic::amdgcn_flat_atomic_fmin:
43414341
case Intrinsic::amdgcn_flat_atomic_fmax:
4342+
case Intrinsic::amdgcn_global_atomic_fadd_v2bf16:
4343+
case Intrinsic::amdgcn_flat_atomic_fadd_v2bf16:
43424344
return getDefaultMappingAllVGPR(MI);
43434345
case Intrinsic::amdgcn_ds_ordered_add:
43444346
case Intrinsic::amdgcn_ds_ordered_swap: {

‎llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -205,9 +205,12 @@ def : SourceOfDivergence<int_amdgcn_global_atomic_fmax>;
205205
def : SourceOfDivergence<int_amdgcn_flat_atomic_fadd>;
206206
def : SourceOfDivergence<int_amdgcn_flat_atomic_fmin>;
207207
def : SourceOfDivergence<int_amdgcn_flat_atomic_fmax>;
208+
def : SourceOfDivergence<int_amdgcn_global_atomic_fadd_v2bf16>;
209+
def : SourceOfDivergence<int_amdgcn_flat_atomic_fadd_v2bf16>;
208210
def : SourceOfDivergence<int_amdgcn_ds_fadd>;
209211
def : SourceOfDivergence<int_amdgcn_ds_fmin>;
210212
def : SourceOfDivergence<int_amdgcn_ds_fmax>;
213+
def : SourceOfDivergence<int_amdgcn_ds_fadd_v2bf16>;
211214
def : SourceOfDivergence<int_amdgcn_buffer_atomic_swap>;
212215
def : SourceOfDivergence<int_amdgcn_buffer_atomic_add>;
213216
def : SourceOfDivergence<int_amdgcn_buffer_atomic_sub>;

‎llvm/lib/Target/AMDGPU/DSInstructions.td‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -462,6 +462,13 @@ let SubtargetPredicate = isGFX90APlus in {
462462
defm DS_ADD_RTN_F64 : DS_1A1D_RET_mc_gfx9<"ds_add_rtn_f64", VReg_64, "ds_add_f64">;
463463
} // End SubtargetPredicate = isGFX90APlus
464464

465+
let SubtargetPredicate = isGFX940Plus in {
466+
defm DS_PK_ADD_F16 : DS_1A1D_NORET_mc_gfx9<"ds_pk_add_f16">;
467+
defm DS_PK_ADD_RTN_F16 : DS_1A1D_RET_mc_gfx9<"ds_pk_add_rtn_f16", VGPR_32, "ds_pk_add_f16">;
468+
defm DS_PK_ADD_BF16 : DS_1A1D_NORET_mc_gfx9<"ds_pk_add_bf16">;
469+
defm DS_PK_ADD_RTN_BF16 : DS_1A1D_RET_mc_gfx9<"ds_pk_add_rtn_bf16", VGPR_32, "ds_pk_add_bf16">;
470+
} // End SubtargetPredicate = isGFX940Plus
471+
465472
defm DS_MSKOR_B32 : DS_1A2D_NORET_mc<"ds_mskor_b32">;
466473
defm DS_CMPST_B32 : DS_1A2D_NORET_mc<"ds_cmpst_b32">;
467474
defm DS_CMPST_F32 : DS_1A2D_NORET_mc<"ds_cmpst_f32">;
@@ -998,6 +1005,14 @@ def : DSAtomicRetPat<DS_ADD_RTN_F64, f64, atomic_load_fadd_local_ret_64>;
9981005
def : DSAtomicRetPat<DS_ADD_F64, f64, atomic_load_fadd_local_noret_64>;
9991006
}
10001007

1008+
let SubtargetPredicate = isGFX940Plus in {
1009+
def : DSAtomicRetPat<DS_PK_ADD_RTN_F16, v2f16, atomic_load_fadd_v2f16_local_32>;
1010+
def : GCNPat <
1011+
(v2i16 (int_amdgcn_ds_fadd_v2bf16 i32:$ptr, v2i16:$src)),
1012+
(DS_PK_ADD_RTN_BF16 VGPR_32:$ptr, VGPR_32:$src, 0, 0)
1013+
>;
1014+
}
1015+
10011016
def : Pat <
10021017
(SIds_ordered_count i32:$value, i16:$offset),
10031018
(DS_ORDERED_COUNT $value, (as_i16imm $offset))
@@ -1410,3 +1425,10 @@ let SubtargetPredicate = isGFX90APlus in {
14101425
def DS_ADD_F64_vi : DS_Real_vi<0x5c, DS_ADD_F64>;
14111426
def DS_ADD_RTN_F64_vi : DS_Real_vi<0x7c, DS_ADD_RTN_F64>;
14121427
} // End SubtargetPredicate = isGFX90APlus
1428+
1429+
let SubtargetPredicate = isGFX940Plus in {
1430+
def DS_PK_ADD_F16_vi : DS_Real_vi<0x17, DS_PK_ADD_F16>;
1431+
def DS_PK_ADD_RTN_F16_vi : DS_Real_vi<0xb7, DS_PK_ADD_RTN_F16>;
1432+
def DS_PK_ADD_BF16_vi : DS_Real_vi<0x18, DS_PK_ADD_BF16>;
1433+
def DS_PK_ADD_RTN_BF16_vi : DS_Real_vi<0xb8, DS_PK_ADD_RTN_BF16>;
1434+
} // End SubtargetPredicate = isGFX940Plus

0 commit comments

Comments
 (0)