Skip to content

Commit 716364e

Browse files
arsenmsrpande
arsenm
and
srpande
authored
AMDGPU: Add support for v_dot2c_f32_bf16 instruction for gfx950 (#117598)
The encoding of v_dot2c_f32_bf16 opcode is same as v_mac_f32 in gfx90a, both from gfx9 series. This required a new decoderNameSpace GFX950_DOT. Co-authored-by: Sirish Pande <[email protected]>
1 parent aa7eb57 commit 716364e

File tree

15 files changed

+372
-15
lines changed

15 files changed

+372
-15
lines changed

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -276,6 +276,7 @@ TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_bf8, "fUiUif", "nc", "dot11-insts")
276276
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_fp8, "fUiUif", "nc", "dot11-insts")
277277
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_fp8, "fUiUif", "nc", "dot11-insts")
278278
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_bf8, "fUiUif", "nc", "dot11-insts")
279+
TARGET_BUILTIN(__builtin_amdgcn_fdot2c_f32_bf16, "fV2yV2yfIb", "nc", "dot13-insts")
279280

280281
//===----------------------------------------------------------------------===//
281282
// GFX10+ only builtins.

‎clang/test/CodeGenOpenCL/amdgpu-features.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -89,7 +89,7 @@
8989
// GFX941: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
9090
// GFX942: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64,+xf32-insts"
9191
// GFX9_4_Generic: "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
92-
// GFX950: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot12-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
92+
// GFX950: "target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+mai-insts,+permlane16-swap,+permlane32-swap,+prng-inst,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
9393
// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
9494
// GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
9595
// GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"

‎clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7,12 +7,14 @@ typedef unsigned int uint;
77
typedef half __attribute__((ext_vector_type(2))) half2;
88
typedef short __attribute__((ext_vector_type(2))) short2;
99
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
10+
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
1011

1112
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
1213
kernel void builtins_amdgcn_dl_insts_err(
1314
global float *fOut, global int *siOut, global uint *uiOut,
1415
global short *sOut, global int *iOut, global half *hOut,
1516
half2 v2hA, half2 v2hB, float fC, half hC,
17+
bfloat2 v2bfbfA, bfloat2 v2bfbfB,
1618
short2 v2ssA, short2 v2ssB, short sC, int siA, int siB, int siC,
1719
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC,
1820
int A, int B, int C) {
@@ -26,6 +28,9 @@ kernel void builtins_amdgcn_dl_insts_err(
2628
fOut[3] = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, false); // expected-error {{'__builtin_amdgcn_fdot2_f32_bf16' needs target feature dot12-insts}}
2729
fOut[4] = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, true); // expected-error {{'__builtin_amdgcn_fdot2_f32_bf16' needs target feature dot12-insts}}
2830

31+
fOut[3] = __builtin_amdgcn_fdot2c_f32_bf16(v2bfbfA, v2bfbfB, fC, false); // expected-error {{'__builtin_amdgcn_fdot2c_f32_bf16' needs target feature dot13-insts}}
32+
fOut[4] = __builtin_amdgcn_fdot2c_f32_bf16(v2bfbfA, v2bfbfB, fC, true); // expected-error {{'__builtin_amdgcn_fdot2c_f32_bf16' needs target feature dot13-insts}}
33+
2934
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, false); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
3035
siOut[1] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, true); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
3136

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl

Lines changed: 28 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@ typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
88
typedef __bf16 __attribute__((ext_vector_type(32))) bfloat32;
99
typedef half __attribute__((ext_vector_type(32))) half32;
1010
typedef short __attribute__((ext_vector_type(2))) short2;
11+
typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
1112
typedef float __attribute__((ext_vector_type(16))) float16;
1213

1314
// CHECK-LABEL: @test_prng_b32(
@@ -216,17 +217,16 @@ void test_ashr_pk_u8_i32(global int* out, uint src0, uint src1, uint src2) {
216217
*out = __builtin_amdgcn_ashr_pk_u8_i32(src0, src1, src2);
217218
}
218219

219-
// CHECK-LABEL: define dso_local void @builtins_amdgcn_dl_insts(
220-
// CHECK-SAME: ptr addrspace(1) noundef [[OUT:%.*]], float noundef [[FC:%.*]], <2 x i16> noundef [[V2SSA:%.*]], <2 x i16> noundef [[V2SSB:%.*]]) #[[ATTR0:[0-9]+]] {
220+
// CHECK-LABEL: @builtins_amdgcn_dl_insts(
221221
// CHECK-NEXT: entry:
222222
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
223223
// CHECK-NEXT: [[FC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
224224
// CHECK-NEXT: [[V2SSA_ADDR:%.*]] = alloca <2 x i16>, align 4, addrspace(5)
225225
// CHECK-NEXT: [[V2SSB_ADDR:%.*]] = alloca <2 x i16>, align 4, addrspace(5)
226-
// CHECK-NEXT: store ptr addrspace(1) [[OUT]], ptr addrspace(5) [[OUT_ADDR]], align 8
227-
// CHECK-NEXT: store float [[FC]], ptr addrspace(5) [[FC_ADDR]], align 4
228-
// CHECK-NEXT: store <2 x i16> [[V2SSA]], ptr addrspace(5) [[V2SSA_ADDR]], align 4
229-
// CHECK-NEXT: store <2 x i16> [[V2SSB]], ptr addrspace(5) [[V2SSB_ADDR]], align 4
226+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
227+
// CHECK-NEXT: store float [[FC:%.*]], ptr addrspace(5) [[FC_ADDR]], align 4
228+
// CHECK-NEXT: store <2 x i16> [[V2SSA:%.*]], ptr addrspace(5) [[V2SSA_ADDR]], align 4
229+
// CHECK-NEXT: store <2 x i16> [[V2SSB:%.*]], ptr addrspace(5) [[V2SSB_ADDR]], align 4
230230
// CHECK-NEXT: [[TMP0:%.*]] = load <2 x i16>, ptr addrspace(5) [[V2SSA_ADDR]], align 4
231231
// CHECK-NEXT: [[TMP1:%.*]] = bitcast <2 x i16> [[TMP0]] to <2 x bfloat>
232232
// CHECK-NEXT: [[TMP2:%.*]] = load <2 x i16>, ptr addrspace(5) [[V2SSB_ADDR]], align 4
@@ -240,3 +240,25 @@ void test_ashr_pk_u8_i32(global int* out, uint src0, uint src1, uint src2) {
240240
void builtins_amdgcn_dl_insts(global float *out, float fC, short2 v2ssA, short2 v2ssB) {
241241
*out = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, false);
242242
}
243+
244+
// CHECK-LABEL: @builtins_amdgcn_dl_dot2c(
245+
// CHECK-NEXT: entry:
246+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
247+
// CHECK-NEXT: [[FC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
248+
// CHECK-NEXT: [[V2SSA_ADDR:%.*]] = alloca <2 x bfloat>, align 4, addrspace(5)
249+
// CHECK-NEXT: [[V2SSB_ADDR:%.*]] = alloca <2 x bfloat>, align 4, addrspace(5)
250+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
251+
// CHECK-NEXT: store float [[FC:%.*]], ptr addrspace(5) [[FC_ADDR]], align 4
252+
// CHECK-NEXT: store <2 x bfloat> [[V2SSA:%.*]], ptr addrspace(5) [[V2SSA_ADDR]], align 4
253+
// CHECK-NEXT: store <2 x bfloat> [[V2SSB:%.*]], ptr addrspace(5) [[V2SSB_ADDR]], align 4
254+
// CHECK-NEXT: [[TMP0:%.*]] = load <2 x bfloat>, ptr addrspace(5) [[V2SSA_ADDR]], align 4
255+
// CHECK-NEXT: [[TMP1:%.*]] = load <2 x bfloat>, ptr addrspace(5) [[V2SSB_ADDR]], align 4
256+
// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(5) [[FC_ADDR]], align 4
257+
// CHECK-NEXT: [[TMP3:%.*]] = call float @llvm.amdgcn.fdot2c.f32.bf16(<2 x bfloat> [[TMP0]], <2 x bfloat> [[TMP1]], float [[TMP2]], i1 false)
258+
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
259+
// CHECK-NEXT: store float [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
260+
// CHECK-NEXT: ret void
261+
//
262+
void builtins_amdgcn_dl_dot2c(global float *out, float fC, bfloat2 v2ssA, bfloat2 v2ssB) {
263+
*out = __builtin_amdgcn_fdot2c_f32_bf16(v2ssA, v2ssB, fC, false);
264+
}

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

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2820,6 +2820,24 @@ def int_amdgcn_fdot2_f32_bf16 :
28202820
[IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<3>>]
28212821
>;
28222822

2823+
// f32 %r = llvm.amdgcn.fdot2c.f32.bf16(v2bf16 %a, v2bf16 %b, f32 %c, i1 %clamp)
2824+
// %r = %a[0] * %b[0] + %a[1] * %b[1] + c
2825+
// TODO: This actually is similar to llvm.amdgcn.fdot2 intrinsics which produces
2826+
// v_dot2c_f32_f16 on gfx940. Maybe we can consolidate these.
2827+
2828+
def int_amdgcn_fdot2c_f32_bf16 :
2829+
ClangBuiltin<"__builtin_amdgcn_fdot2c_f32_bf16">,
2830+
DefaultAttrsIntrinsic<
2831+
[llvm_float_ty], // %r
2832+
[
2833+
llvm_v2bf16_ty, // %a
2834+
llvm_v2bf16_ty, // %b
2835+
llvm_float_ty, // %c
2836+
llvm_i1_ty // %clamp
2837+
],
2838+
[IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<3>>]
2839+
>;
2840+
28232841
// i32 %r = llvm.amdgcn.sdot2(v2i16 %a, v2i16 %b, i32 %c, i1 %clamp)
28242842
// %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
28252843
def int_amdgcn_sdot2 :

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

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -753,6 +753,13 @@ def FeatureDot12Insts : SubtargetFeature<"dot12-insts",
753753
"Has v_dot2_f32_bf16 instructions"
754754
>;
755755

756+
def FeatureDot13Insts : SubtargetFeature<"dot13-insts",
757+
"HasDot13Insts",
758+
"true",
759+
"Has v_dot2c_f32_bf16 instructions"
760+
>;
761+
762+
756763
def FeatureMAIInsts : SubtargetFeature<"mai-insts",
757764
"HasMAIInsts",
758765
"true",
@@ -1585,7 +1592,8 @@ def FeatureISAVersion9_5_Common : FeatureSet<
15851592
FeatureBF8ConversionScaleInsts,
15861593
FeatureFP4ConversionScaleInsts,
15871594
FeatureFP6BF6ConversionScaleInsts,
1588-
FeatureDot12Insts
1595+
FeatureDot12Insts,
1596+
FeatureDot13Insts
15891597
])>;
15901598

15911599
def FeatureISAVersion9_4_0 : FeatureSet<
@@ -2373,6 +2381,9 @@ def HasDot11Insts : Predicate<"Subtarget->hasDot11Insts()">,
23732381
def HasDot12Insts : Predicate<"Subtarget->hasDot12Insts()">,
23742382
AssemblerPredicate<(all_of FeatureDot12Insts)>;
23752383

2384+
def HasDot13Insts : Predicate<"Subtarget->hasDot13Insts()">,
2385+
AssemblerPredicate<(all_of FeatureDot13Insts)>;
2386+
23762387
def HasGetWaveIdInst : Predicate<"Subtarget->hasGetWaveIdInst()">,
23772388
AssemblerPredicate<(all_of FeatureGetWaveIdInst)>;
23782389

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4528,6 +4528,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
45284528
case Intrinsic::amdgcn_fdot2_bf16_bf16:
45294529
case Intrinsic::amdgcn_fdot2_f16_f16:
45304530
case Intrinsic::amdgcn_fdot2_f32_bf16:
4531+
case Intrinsic::amdgcn_fdot2c_f32_bf16:
45314532
case Intrinsic::amdgcn_sudot4:
45324533
case Intrinsic::amdgcn_sudot8:
45334534
case Intrinsic::amdgcn_dot4_f32_fp8_bf8:

‎llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -566,6 +566,10 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size,
566566
tryDecodeInst(DecoderTableGFX80_UNPACKED64, MI, QW, Address, CS))
567567
break;
568568

569+
if (STI.hasFeature(AMDGPU::FeatureGFX950Insts) &&
570+
tryDecodeInst(DecoderTableGFX95064, MI, QW, Address, CS))
571+
break;
572+
569573
// Some GFX9 subtargets repurposed the v_mad_mix_f32, v_mad_mixlo_f16 and
570574
// v_mad_mixhi_f16 for FMA variants. Try to decode using this special
571575
// table first so we print the correct name.
@@ -627,6 +631,10 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size,
627631
if (isGFX9() && tryDecodeInst(DecoderTableGFX932, MI, DW, Address, CS))
628632
break;
629633

634+
if (STI.hasFeature(AMDGPU::FeatureGFX950Insts) &&
635+
tryDecodeInst(DecoderTableGFX95032, MI, DW, Address, CS))
636+
break;
637+
630638
if (STI.hasFeature(AMDGPU::FeatureGFX90AInsts) &&
631639
tryDecodeInst(DecoderTableGFX90A32, MI, DW, Address, CS))
632640
break;

‎llvm/lib/Target/AMDGPU/GCNSubtarget.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -157,6 +157,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
157157
bool HasDot10Insts = false;
158158
bool HasDot11Insts = false;
159159
bool HasDot12Insts = false;
160+
bool HasDot13Insts = false;
160161
bool HasMAIInsts = false;
161162
bool HasFP8Insts = false;
162163
bool HasFP8ConversionInsts = false;
@@ -830,6 +831,10 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
830831
return HasDot12Insts;
831832
}
832833

834+
bool hasDot13Insts() const {
835+
return HasDot13Insts;
836+
}
837+
833838
bool hasMAIInsts() const {
834839
return HasMAIInsts;
835840
}

‎llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -591,6 +591,7 @@ bool isMAC(unsigned Opc) {
591591
Opc == AMDGPU::V_FMAC_F16_fake16_e64_gfx11 ||
592592
Opc == AMDGPU::V_FMAC_F16_fake16_e64_gfx12 ||
593593
Opc == AMDGPU::V_DOT2C_F32_F16_e64_vi ||
594+
Opc == AMDGPU::V_DOT2C_F32_BF16_e64_vi ||
594595
Opc == AMDGPU::V_DOT2C_I32_I16_e64_vi ||
595596
Opc == AMDGPU::V_DOT4C_I32_I8_e64_vi ||
596597
Opc == AMDGPU::V_DOT8C_I32_I4_e64_vi;

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

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -567,6 +567,12 @@ def VOP_DOT_ACC_F32_V2F16 : VOP_DOT_ACC<f32, v2f16> {
567567
let HasClamp = 1;
568568
}
569569

570+
def VOP_DOT_ACC_F32_V2BF16 : VOP_DOT_ACC<f32, v2bf16> {
571+
let Src0ModDPP = FPVRegInputMods;
572+
let Src1ModDPP = FPVRegInputMods;
573+
let HasClamp = 1;
574+
}
575+
570576
def VOP_DOT_ACC_I32_I32 : VOP_DOT_ACC<i32, i32> {
571577
let HasExtVOP3DPP = 0;
572578
let HasSrc0Mods = 1;
@@ -1182,6 +1188,9 @@ let Constraints = "$vdst = $src2",
11821188
defm V_DOT2C_I32_I16 : VOP2Inst<"v_dot2c_i32_i16", VOP_DOT_ACC_I32_I32>;
11831189
let SubtargetPredicate = HasDot3Insts in
11841190
defm V_DOT8C_I32_I4 : VOP2Inst<"v_dot8c_i32_i4", VOP_DOT_ACC_I32_I32>;
1191+
1192+
let SubtargetPredicate = HasDot13Insts in
1193+
defm V_DOT2C_F32_BF16 : VOP2Inst<"v_dot2c_f32_bf16", VOP_DOT_ACC_F32_V2BF16>;
11851194
}
11861195

11871196
let AddedComplexity = 30 in {
@@ -1191,6 +1200,12 @@ let AddedComplexity = 30 in {
11911200
> {
11921201
let SubtargetPredicate = HasDot5Insts;
11931202
}
1203+
def : GCNPat<
1204+
(f32 (int_amdgcn_fdot2_f32_bf16 v2bf16:$src0, v2bf16:$src1, f32:$src2, (i1 DSTCLAMP.NONE))),
1205+
(f32 (V_DOT2C_F32_BF16_e32 $src0, $src1, $src2))
1206+
> {
1207+
let SubtargetPredicate = HasDot13Insts;
1208+
}
11941209
def : GCNPat<
11951210
(i32 (int_amdgcn_sdot4 i32:$src0, i32:$src1, i32:$src2, (i1 DSTCLAMP.NONE))),
11961211
(i32 (V_DOT4C_I32_I8_e32 $src0, $src1, $src2))
@@ -2670,3 +2685,8 @@ let SubtargetPredicate = HasDot3Insts in {
26702685
let DecoderNamespace = "GFX10_B" in
26712686
defm V_DOT8C_I32_I4 : VOP2_Real_DOT_ACC_gfx10<0x02>;
26722687
}
2688+
2689+
let OtherPredicates = [HasDot13Insts] in {
2690+
let DecoderNamespace = "GFX950" in
2691+
defm V_DOT2C_F32_BF16 : VOP2_Real_DOT_ACC_gfx9<0x16>;
2692+
}

‎llvm/lib/TargetParser/TargetParser.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -478,6 +478,7 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
478478
Features["permlane32-swap"] = true;
479479
Features["ashr-pk-insts"] = true;
480480
Features["dot12-insts"] = true;
481+
Features["dot13-insts"] = true;
481482
Features["gfx950-insts"] = true;
482483
[[fallthrough]];
483484
case GK_GFX942:

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fdot2.f32.bf16.ll

Lines changed: 5 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -52,7 +52,6 @@ define amdgpu_kernel void @test_llvm_amdgcn_fdot2_f32_bf16_clamp(
5252
; GFX950-ISEL-NEXT: s_nop 1
5353
; GFX950-ISEL-NEXT: global_store_dword v1, v0, s[8:9]
5454
; GFX950-ISEL-NEXT: s_endpgm
55-
5655
ptr addrspace(1) %r,
5756
ptr addrspace(1) %a,
5857
ptr addrspace(1) %b,
@@ -93,9 +92,9 @@ define amdgpu_kernel void @test_llvm_amdgcn_fdot2_f32_bf16_no_clamp(
9392
; GFX950-NEXT: s_waitcnt lgkmcnt(0)
9493
; GFX950-NEXT: v_mov_b32_e32 v1, s0
9594
; GFX950-NEXT: v_mov_b32_e32 v2, s1
96-
; GFX950-NEXT: v_dot2_f32_bf16 v1, s2, v1, v2
95+
; GFX950-NEXT: v_dot2c_f32_bf16_e32 v2, s2, v1
9796
; GFX950-NEXT: s_nop 2
98-
; GFX950-NEXT: global_store_dword v0, v1, s[8:9]
97+
; GFX950-NEXT: global_store_dword v0, v2, s[8:9]
9998
; GFX950-NEXT: s_endpgm
10099
;
101100
; GFX950-ISEL-LABEL: test_llvm_amdgcn_fdot2_f32_bf16_no_clamp:
@@ -108,12 +107,11 @@ define amdgpu_kernel void @test_llvm_amdgcn_fdot2_f32_bf16_no_clamp(
108107
; GFX950-ISEL-NEXT: s_waitcnt lgkmcnt(0)
109108
; GFX950-ISEL-NEXT: v_mov_b32_e32 v0, s0
110109
; GFX950-ISEL-NEXT: v_mov_b32_e32 v1, s1
111-
; GFX950-ISEL-NEXT: v_dot2_f32_bf16 v0, s2, v0, v1
112-
; GFX950-ISEL-NEXT: v_mov_b32_e32 v1, 0
110+
; GFX950-ISEL-NEXT: v_dot2c_f32_bf16_e32 v1, s2, v0
111+
; GFX950-ISEL-NEXT: v_mov_b32_e32 v0, 0
113112
; GFX950-ISEL-NEXT: s_nop 1
114-
; GFX950-ISEL-NEXT: global_store_dword v1, v0, s[8:9]
113+
; GFX950-ISEL-NEXT: global_store_dword v0, v1, s[8:9]
115114
; GFX950-ISEL-NEXT: s_endpgm
116-
117115
ptr addrspace(1) %r,
118116
ptr addrspace(1) %a,
119117
ptr addrspace(1) %b,

0 commit comments

Comments
 (0)