[clang] 716364e - AMDGPU: Add support for v_dot2c_f32_bf16 instruction for gfx950 (#117598)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 25 19:51:04 PST 2024
Author: Matt Arsenault
Date: 2024-11-25T19:51:01-08:00
New Revision: 716364ebd6649aeca8658680ebb8b0424d028006
URL: https://github.com/llvm/llvm-project/commit/716364ebd6649aeca8658680ebb8b0424d028006
DIFF: https://github.com/llvm/llvm-project/commit/716364ebd6649aeca8658680ebb8b0424d028006.diff
LOG: 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 <Sirish.Pande at amd.com>
Added:
llvm/test/MC/AMDGPU/gfx950_xdlops.s
llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_xdlops.txt
Modified:
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/test/CodeGenOpenCL/amdgpu-features.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/lib/Target/AMDGPU/AMDGPU.td
llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
llvm/lib/Target/AMDGPU/GCNSubtarget.h
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
llvm/lib/Target/AMDGPU/VOP2Instructions.td
llvm/lib/TargetParser/TargetParser.cpp
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fdot2.f32.bf16.ll
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 7d0019eead96b6..49304d12d6d70d 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -276,6 +276,7 @@ TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_bf8, "fUiUif", "nc", "dot11-insts")
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_fp8, "fUiUif", "nc", "dot11-insts")
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_fp8, "fUiUif", "nc", "dot11-insts")
TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_bf8, "fUiUif", "nc", "dot11-insts")
+TARGET_BUILTIN(__builtin_amdgcn_fdot2c_f32_bf16, "fV2yV2yfIb", "nc", "dot13-insts")
//===----------------------------------------------------------------------===//
// GFX10+ only builtins.
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl
index 0b698035ee54c7..1e2921160d28f2 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl
@@ -89,7 +89,7 @@
// 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"
// 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"
// 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"
-// 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"
+// 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"
// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
// 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"
// 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"
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
index f409d2a110d753..b8c46039dac530 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
@@ -7,12 +7,14 @@ typedef unsigned int uint;
typedef half __attribute__((ext_vector_type(2))) half2;
typedef short __attribute__((ext_vector_type(2))) short2;
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
+typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
kernel void builtins_amdgcn_dl_insts_err(
global float *fOut, global int *siOut, global uint *uiOut,
global short *sOut, global int *iOut, global half *hOut,
half2 v2hA, half2 v2hB, float fC, half hC,
+ bfloat2 v2bfbfA, bfloat2 v2bfbfB,
short2 v2ssA, short2 v2ssB, short sC, int siA, int siB, int siC,
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC,
int A, int B, int C) {
@@ -26,6 +28,9 @@ kernel void builtins_amdgcn_dl_insts_err(
fOut[3] = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, false); // expected-error {{'__builtin_amdgcn_fdot2_f32_bf16' needs target feature dot12-insts}}
fOut[4] = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, true); // expected-error {{'__builtin_amdgcn_fdot2_f32_bf16' needs target feature dot12-insts}}
+ fOut[3] = __builtin_amdgcn_fdot2c_f32_bf16(v2bfbfA, v2bfbfB, fC, false); // expected-error {{'__builtin_amdgcn_fdot2c_f32_bf16' needs target feature dot13-insts}}
+ fOut[4] = __builtin_amdgcn_fdot2c_f32_bf16(v2bfbfA, v2bfbfB, fC, true); // expected-error {{'__builtin_amdgcn_fdot2c_f32_bf16' needs target feature dot13-insts}}
+
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, false); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
siOut[1] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, true); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dot2-insts}}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
index 390d9a55e9b16d..d2125e90bc2c89 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx950.cl
@@ -8,6 +8,7 @@ typedef unsigned int __attribute__((ext_vector_type(6))) uint6;
typedef __bf16 __attribute__((ext_vector_type(32))) bfloat32;
typedef half __attribute__((ext_vector_type(32))) half32;
typedef short __attribute__((ext_vector_type(2))) short2;
+typedef __bf16 __attribute__((ext_vector_type(2))) bfloat2;
typedef float __attribute__((ext_vector_type(16))) float16;
// CHECK-LABEL: @test_prng_b32(
@@ -216,17 +217,16 @@ void test_ashr_pk_u8_i32(global int* out, uint src0, uint src1, uint src2) {
*out = __builtin_amdgcn_ashr_pk_u8_i32(src0, src1, src2);
}
-// CHECK-LABEL: define dso_local void @builtins_amdgcn_dl_insts(
-// CHECK-SAME: ptr addrspace(1) noundef [[OUT:%.*]], float noundef [[FC:%.*]], <2 x i16> noundef [[V2SSA:%.*]], <2 x i16> noundef [[V2SSB:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-LABEL: @builtins_amdgcn_dl_insts(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
// CHECK-NEXT: [[FC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
// CHECK-NEXT: [[V2SSA_ADDR:%.*]] = alloca <2 x i16>, align 4, addrspace(5)
// CHECK-NEXT: [[V2SSB_ADDR:%.*]] = alloca <2 x i16>, align 4, addrspace(5)
-// CHECK-NEXT: store ptr addrspace(1) [[OUT]], ptr addrspace(5) [[OUT_ADDR]], align 8
-// CHECK-NEXT: store float [[FC]], ptr addrspace(5) [[FC_ADDR]], align 4
-// CHECK-NEXT: store <2 x i16> [[V2SSA]], ptr addrspace(5) [[V2SSA_ADDR]], align 4
-// CHECK-NEXT: store <2 x i16> [[V2SSB]], ptr addrspace(5) [[V2SSB_ADDR]], align 4
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store float [[FC:%.*]], ptr addrspace(5) [[FC_ADDR]], align 4
+// CHECK-NEXT: store <2 x i16> [[V2SSA:%.*]], ptr addrspace(5) [[V2SSA_ADDR]], align 4
+// CHECK-NEXT: store <2 x i16> [[V2SSB:%.*]], ptr addrspace(5) [[V2SSB_ADDR]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load <2 x i16>, ptr addrspace(5) [[V2SSA_ADDR]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = bitcast <2 x i16> [[TMP0]] to <2 x bfloat>
// 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) {
void builtins_amdgcn_dl_insts(global float *out, float fC, short2 v2ssA, short2 v2ssB) {
*out = __builtin_amdgcn_fdot2_f32_bf16(v2ssA, v2ssB, fC, false);
}
+
+// CHECK-LABEL: @builtins_amdgcn_dl_dot2c(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
+// CHECK-NEXT: [[FC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// CHECK-NEXT: [[V2SSA_ADDR:%.*]] = alloca <2 x bfloat>, align 4, addrspace(5)
+// CHECK-NEXT: [[V2SSB_ADDR:%.*]] = alloca <2 x bfloat>, align 4, addrspace(5)
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store float [[FC:%.*]], ptr addrspace(5) [[FC_ADDR]], align 4
+// CHECK-NEXT: store <2 x bfloat> [[V2SSA:%.*]], ptr addrspace(5) [[V2SSA_ADDR]], align 4
+// CHECK-NEXT: store <2 x bfloat> [[V2SSB:%.*]], ptr addrspace(5) [[V2SSB_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load <2 x bfloat>, ptr addrspace(5) [[V2SSA_ADDR]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load <2 x bfloat>, ptr addrspace(5) [[V2SSB_ADDR]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(5) [[FC_ADDR]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = call float @llvm.amdgcn.fdot2c.f32.bf16(<2 x bfloat> [[TMP0]], <2 x bfloat> [[TMP1]], float [[TMP2]], i1 false)
+// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[OUT_ADDR]], align 8
+// CHECK-NEXT: store float [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
+// CHECK-NEXT: ret void
+//
+void builtins_amdgcn_dl_dot2c(global float *out, float fC, bfloat2 v2ssA, bfloat2 v2ssB) {
+ *out = __builtin_amdgcn_fdot2c_f32_bf16(v2ssA, v2ssB, fC, false);
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index b8e07cc799b42e..c35c1635a7b6b5 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2820,6 +2820,24 @@ def int_amdgcn_fdot2_f32_bf16 :
[IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<3>>]
>;
+// f32 %r = llvm.amdgcn.fdot2c.f32.bf16(v2bf16 %a, v2bf16 %b, f32 %c, i1 %clamp)
+// %r = %a[0] * %b[0] + %a[1] * %b[1] + c
+// TODO: This actually is similar to llvm.amdgcn.fdot2 intrinsics which produces
+// v_dot2c_f32_f16 on gfx940. Maybe we can consolidate these.
+
+def int_amdgcn_fdot2c_f32_bf16 :
+ ClangBuiltin<"__builtin_amdgcn_fdot2c_f32_bf16">,
+ DefaultAttrsIntrinsic<
+ [llvm_float_ty], // %r
+ [
+ llvm_v2bf16_ty, // %a
+ llvm_v2bf16_ty, // %b
+ llvm_float_ty, // %c
+ llvm_i1_ty // %clamp
+ ],
+ [IntrNoMem, IntrSpeculatable, ImmArg<ArgIndex<3>>]
+ >;
+
// i32 %r = llvm.amdgcn.sdot2(v2i16 %a, v2i16 %b, i32 %c, i1 %clamp)
// %r = %a[0] * %b[0] + %a[1] * %b[1] + %c
def int_amdgcn_sdot2 :
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 0a53f661f33657..be213431ea8f1f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -753,6 +753,13 @@ def FeatureDot12Insts : SubtargetFeature<"dot12-insts",
"Has v_dot2_f32_bf16 instructions"
>;
+def FeatureDot13Insts : SubtargetFeature<"dot13-insts",
+ "HasDot13Insts",
+ "true",
+ "Has v_dot2c_f32_bf16 instructions"
+>;
+
+
def FeatureMAIInsts : SubtargetFeature<"mai-insts",
"HasMAIInsts",
"true",
@@ -1585,7 +1592,8 @@ def FeatureISAVersion9_5_Common : FeatureSet<
FeatureBF8ConversionScaleInsts,
FeatureFP4ConversionScaleInsts,
FeatureFP6BF6ConversionScaleInsts,
- FeatureDot12Insts
+ FeatureDot12Insts,
+ FeatureDot13Insts
])>;
def FeatureISAVersion9_4_0 : FeatureSet<
@@ -2373,6 +2381,9 @@ def HasDot11Insts : Predicate<"Subtarget->hasDot11Insts()">,
def HasDot12Insts : Predicate<"Subtarget->hasDot12Insts()">,
AssemblerPredicate<(all_of FeatureDot12Insts)>;
+def HasDot13Insts : Predicate<"Subtarget->hasDot13Insts()">,
+ AssemblerPredicate<(all_of FeatureDot13Insts)>;
+
def HasGetWaveIdInst : Predicate<"Subtarget->hasGetWaveIdInst()">,
AssemblerPredicate<(all_of FeatureGetWaveIdInst)>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index 7651c84eb0ae65..b06bd4e334614f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4528,6 +4528,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
case Intrinsic::amdgcn_fdot2_bf16_bf16:
case Intrinsic::amdgcn_fdot2_f16_f16:
case Intrinsic::amdgcn_fdot2_f32_bf16:
+ case Intrinsic::amdgcn_fdot2c_f32_bf16:
case Intrinsic::amdgcn_sudot4:
case Intrinsic::amdgcn_sudot8:
case Intrinsic::amdgcn_dot4_f32_fp8_bf8:
diff --git a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
index 983a10027b20f4..5908351805721f 100644
--- a/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
+++ b/llvm/lib/Target/AMDGPU/Disassembler/AMDGPUDisassembler.cpp
@@ -566,6 +566,10 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size,
tryDecodeInst(DecoderTableGFX80_UNPACKED64, MI, QW, Address, CS))
break;
+ if (STI.hasFeature(AMDGPU::FeatureGFX950Insts) &&
+ tryDecodeInst(DecoderTableGFX95064, MI, QW, Address, CS))
+ break;
+
// Some GFX9 subtargets repurposed the v_mad_mix_f32, v_mad_mixlo_f16 and
// v_mad_mixhi_f16 for FMA variants. Try to decode using this special
// table first so we print the correct name.
@@ -627,6 +631,10 @@ DecodeStatus AMDGPUDisassembler::getInstruction(MCInst &MI, uint64_t &Size,
if (isGFX9() && tryDecodeInst(DecoderTableGFX932, MI, DW, Address, CS))
break;
+ if (STI.hasFeature(AMDGPU::FeatureGFX950Insts) &&
+ tryDecodeInst(DecoderTableGFX95032, MI, DW, Address, CS))
+ break;
+
if (STI.hasFeature(AMDGPU::FeatureGFX90AInsts) &&
tryDecodeInst(DecoderTableGFX90A32, MI, DW, Address, CS))
break;
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index 7e994b84426bf0..cdc5e1a66afa2c 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -157,6 +157,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
bool HasDot10Insts = false;
bool HasDot11Insts = false;
bool HasDot12Insts = false;
+ bool HasDot13Insts = false;
bool HasMAIInsts = false;
bool HasFP8Insts = false;
bool HasFP8ConversionInsts = false;
@@ -830,6 +831,10 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
return HasDot12Insts;
}
+ bool hasDot13Insts() const {
+ return HasDot13Insts;
+ }
+
bool hasMAIInsts() const {
return HasMAIInsts;
}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index c7e4659b15d299..b233e898589396 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -591,6 +591,7 @@ bool isMAC(unsigned Opc) {
Opc == AMDGPU::V_FMAC_F16_fake16_e64_gfx11 ||
Opc == AMDGPU::V_FMAC_F16_fake16_e64_gfx12 ||
Opc == AMDGPU::V_DOT2C_F32_F16_e64_vi ||
+ Opc == AMDGPU::V_DOT2C_F32_BF16_e64_vi ||
Opc == AMDGPU::V_DOT2C_I32_I16_e64_vi ||
Opc == AMDGPU::V_DOT4C_I32_I8_e64_vi ||
Opc == AMDGPU::V_DOT8C_I32_I4_e64_vi;
diff --git a/llvm/lib/Target/AMDGPU/VOP2Instructions.td b/llvm/lib/Target/AMDGPU/VOP2Instructions.td
index 103575dc351f2b..128c7756191181 100644
--- a/llvm/lib/Target/AMDGPU/VOP2Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP2Instructions.td
@@ -567,6 +567,12 @@ def VOP_DOT_ACC_F32_V2F16 : VOP_DOT_ACC<f32, v2f16> {
let HasClamp = 1;
}
+def VOP_DOT_ACC_F32_V2BF16 : VOP_DOT_ACC<f32, v2bf16> {
+ let Src0ModDPP = FPVRegInputMods;
+ let Src1ModDPP = FPVRegInputMods;
+ let HasClamp = 1;
+}
+
def VOP_DOT_ACC_I32_I32 : VOP_DOT_ACC<i32, i32> {
let HasExtVOP3DPP = 0;
let HasSrc0Mods = 1;
@@ -1182,6 +1188,9 @@ let Constraints = "$vdst = $src2",
defm V_DOT2C_I32_I16 : VOP2Inst<"v_dot2c_i32_i16", VOP_DOT_ACC_I32_I32>;
let SubtargetPredicate = HasDot3Insts in
defm V_DOT8C_I32_I4 : VOP2Inst<"v_dot8c_i32_i4", VOP_DOT_ACC_I32_I32>;
+
+ let SubtargetPredicate = HasDot13Insts in
+ defm V_DOT2C_F32_BF16 : VOP2Inst<"v_dot2c_f32_bf16", VOP_DOT_ACC_F32_V2BF16>;
}
let AddedComplexity = 30 in {
@@ -1191,6 +1200,12 @@ let AddedComplexity = 30 in {
> {
let SubtargetPredicate = HasDot5Insts;
}
+ def : GCNPat<
+ (f32 (int_amdgcn_fdot2_f32_bf16 v2bf16:$src0, v2bf16:$src1, f32:$src2, (i1 DSTCLAMP.NONE))),
+ (f32 (V_DOT2C_F32_BF16_e32 $src0, $src1, $src2))
+ > {
+ let SubtargetPredicate = HasDot13Insts;
+ }
def : GCNPat<
(i32 (int_amdgcn_sdot4 i32:$src0, i32:$src1, i32:$src2, (i1 DSTCLAMP.NONE))),
(i32 (V_DOT4C_I32_I8_e32 $src0, $src1, $src2))
@@ -2670,3 +2685,8 @@ let SubtargetPredicate = HasDot3Insts in {
let DecoderNamespace = "GFX10_B" in
defm V_DOT8C_I32_I4 : VOP2_Real_DOT_ACC_gfx10<0x02>;
}
+
+let OtherPredicates = [HasDot13Insts] in {
+ let DecoderNamespace = "GFX950" in
+ defm V_DOT2C_F32_BF16 : VOP2_Real_DOT_ACC_gfx9<0x16>;
+}
diff --git a/llvm/lib/TargetParser/TargetParser.cpp b/llvm/lib/TargetParser/TargetParser.cpp
index 110b94a57b802c..8b445b1aaa0548 100644
--- a/llvm/lib/TargetParser/TargetParser.cpp
+++ b/llvm/lib/TargetParser/TargetParser.cpp
@@ -478,6 +478,7 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
Features["permlane32-swap"] = true;
Features["ashr-pk-insts"] = true;
Features["dot12-insts"] = true;
+ Features["dot13-insts"] = true;
Features["gfx950-insts"] = true;
[[fallthrough]];
case GK_GFX942:
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fdot2.f32.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fdot2.f32.bf16.ll
index fca418e3f82002..42acf089e86488 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fdot2.f32.bf16.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.fdot2.f32.bf16.ll
@@ -52,7 +52,6 @@ define amdgpu_kernel void @test_llvm_amdgcn_fdot2_f32_bf16_clamp(
; GFX950-ISEL-NEXT: s_nop 1
; GFX950-ISEL-NEXT: global_store_dword v1, v0, s[8:9]
; GFX950-ISEL-NEXT: s_endpgm
-
ptr addrspace(1) %r,
ptr addrspace(1) %a,
ptr addrspace(1) %b,
@@ -93,9 +92,9 @@ define amdgpu_kernel void @test_llvm_amdgcn_fdot2_f32_bf16_no_clamp(
; GFX950-NEXT: s_waitcnt lgkmcnt(0)
; GFX950-NEXT: v_mov_b32_e32 v1, s0
; GFX950-NEXT: v_mov_b32_e32 v2, s1
-; GFX950-NEXT: v_dot2_f32_bf16 v1, s2, v1, v2
+; GFX950-NEXT: v_dot2c_f32_bf16_e32 v2, s2, v1
; GFX950-NEXT: s_nop 2
-; GFX950-NEXT: global_store_dword v0, v1, s[8:9]
+; GFX950-NEXT: global_store_dword v0, v2, s[8:9]
; GFX950-NEXT: s_endpgm
;
; 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(
; GFX950-ISEL-NEXT: s_waitcnt lgkmcnt(0)
; GFX950-ISEL-NEXT: v_mov_b32_e32 v0, s0
; GFX950-ISEL-NEXT: v_mov_b32_e32 v1, s1
-; GFX950-ISEL-NEXT: v_dot2_f32_bf16 v0, s2, v0, v1
-; GFX950-ISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX950-ISEL-NEXT: v_dot2c_f32_bf16_e32 v1, s2, v0
+; GFX950-ISEL-NEXT: v_mov_b32_e32 v0, 0
; GFX950-ISEL-NEXT: s_nop 1
-; GFX950-ISEL-NEXT: global_store_dword v1, v0, s[8:9]
+; GFX950-ISEL-NEXT: global_store_dword v0, v1, s[8:9]
; GFX950-ISEL-NEXT: s_endpgm
-
ptr addrspace(1) %r,
ptr addrspace(1) %a,
ptr addrspace(1) %b,
diff --git a/llvm/test/MC/AMDGPU/gfx950_xdlops.s b/llvm/test/MC/AMDGPU/gfx950_xdlops.s
new file mode 100644
index 00000000000000..2ca131c9c0bf4f
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/gfx950_xdlops.s
@@ -0,0 +1,133 @@
+// RUN: llvm-mc -triple=amdgcn -mcpu=gfx950 -show-encoding %s | FileCheck %s
+
+// CHECK: encoding: [0x01,0x05,0x0a,0x2c]
+v_dot2c_f32_bf16 v5, v1, v2
+
+// CHECK: encoding: [0x01,0x05,0xfe,0x2d]
+v_dot2c_f32_bf16 v255, v1, v2
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0xfe,0x2d,0x01,0xe4,0x00,0x00]
+v_dot2c_f32_bf16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0xff,0xe4,0x00,0x00]
+v_dot2c_f32_bf16_dpp v5, v255, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0xfe,0x0b,0x2c,0x01,0xe4,0x00,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v255 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x1b,0x00,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x40,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_mirror row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x41,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_half_mirror row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x42,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_bcast:15 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x43,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_bcast:31 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x30,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 wave_shl:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x34,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 wave_rol:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x38,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 wave_shr:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x3c,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 wave_ror:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x01,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_shl:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x0f,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_shl:15 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x11,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_shr:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x1f,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_shr:15 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x21,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_ror:1 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x2f,0x01,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 row_ror:15 row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x10]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x1 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x30]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x3 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0xf0]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0xf0]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x01]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x1
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x03]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x3
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x0f]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0xf
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x0f]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0xf
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x08,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 bound_ctrl:1
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x10,0x00]
+v_dot2c_f32_bf16_dpp v5, -v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x20,0x00]
+v_dot2c_f32_bf16_dpp v5, |v1|, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x40,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, -v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x80,0x00]
+v_dot2c_f32_bf16_dpp v5, v1, |v2| quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0
+
+// CHECK: encoding: [0x05,0x00,0x16,0xd1,0x01,0xfb,0x01,0x00]
+v_dot2c_f32_bf16_e64 v5, v1, src_scc
+
+// CHECK: encoding: [0x05,0x00,0x16,0xd1,0xff,0xf9,0x01,0x00]
+v_dot2c_f32_bf16_e64 v5, v255, src_execz
+
+// CHECK: encoding: [0x05,0x00,0x16,0xd1,0x65,0xca,0x00,0x00]
+v_dot2c_f32_bf16_e64 v5, s101, s101
+
+// CHECK: encoding: [0x05,0x00,0x16,0xd1,0xc1,0xcc,0x00,0x00]
+v_dot2c_f32_bf16_e64 v5, -1, flat_scratch_lo
+
+// CHECK: encoding: [0x05,0x02,0x16,0xd1,0xf0,0xce,0x00,0x40]
+v_dot2c_f32_bf16_e64 v5, 0.5, -|flat_scratch_hi|
+
+// CHECK: encoding: [0x05,0x00,0x16,0xd1,0xfc,0xe0,0x01,0x10]
+v_dot2c_f32_bf16_e64 v5, src_execz, 0.5 mul:4
+
+// CHECK: encoding: [0xff,0x81,0x16,0xd1,0xfd,0x82,0x01,0x38]
+v_dot2c_f32_bf16_e64 v255, -|src_scc|, -1 clamp div:2
+
+// CHECK: encoding: [0x8a,0x04,0x0a,0x2c]
+v_dot2c_f32_bf16_e32 v5, 10, v2 ; encoding: [0x8a,0x04,0x0a,0x2c]
+
+// CHECK: encoding: [0xff,0x04,0x0a,0x2c,0x64,0x00,0x00,0x00]
+v_dot2c_f32_bf16_e32 v5, 100, v2 ; encoding: [0xff,0x04,0x0a,0x2c,0x64,0x00,0x00,0x00]
+
+// CHECK: encoding: [0xff,0x04,0x0a,0x2c,0x22,0x41,0x00,0x00]
+v_dot2c_f32_bf16_e32 v5, 10.1, v2 ; encoding: [0xff,0x04,0x0a,0x2c,0x22,0x41,0x00,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_xdlops.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_xdlops.txt
new file mode 100644
index 00000000000000..53b0bcb0aa1ae7
--- /dev/null
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx950_dasm_xdlops.txt
@@ -0,0 +1,133 @@
+# RUN: llvm-mc -arch=amdgcn -mcpu=gfx950 -disassemble -show-encoding < %s | FileCheck -check-prefix=GFX950 %s
+
+# GFX950: v_dot2c_f32_bf16_e32 v5, v1, v2 ; encoding: [0x01,0x05,0x0a,0x2c]
+0x01,0x05,0x0a,0x2c
+
+# GFX950: v_dot2c_f32_bf16_e32 v255, v1, v2 ; encoding: [0x01,0x05,0xfe,0x2d]
+0x01,0x05,0xfe,0x2d
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v255, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0xfe,0x2d,0x01,0xe4,0x00,0x00]
+0xfa,0x04,0xfe,0x2d,0x01,0xe4,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v255, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0xff,0xe4,0x00,0x00]
+0xfa,0x04,0x0a,0x2c,0xff,0xe4,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v255 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0xfe,0x0b,0x2c,0x01,0xe4,0x00,0x00]
+0xfa,0xfe,0x0b,0x2c,0x01,0xe4,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[3,2,1,0] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x1b,0x00,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x1b,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_mirror row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x40,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x40,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_half_mirror row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x41,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x41,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_bcast:15 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x42,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x42,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_bcast:31 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x43,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x43,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 wave_shl:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x30,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x30,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 wave_rol:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x34,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x34,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 wave_shr:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x38,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x38,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 wave_ror:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x3c,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x3c,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_shl:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x01,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x01,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_shl:15 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x0f,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x0f,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_shr:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x11,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x11,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_shr:15 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x1f,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x1f,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_ror:1 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x21,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x21,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 row_ror:15 row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0x2f,0x01,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0x2f,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x1 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x10]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x10
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x3 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x30]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x30
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0xf0]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0xf0
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0xf0]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0xf0
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x01]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x01
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x3 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x03]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x03
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0xf ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x0f]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x0f
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0xf ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x0f]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x00,0x0f
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 bound_ctrl:1 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x08,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x08,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, -v1, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x10,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x10,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, |v1|, v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x20,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x20,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, -v2 quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x40,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x40,0x00
+
+# GFX950: v_dot2c_f32_bf16_dpp v5, v1, |v2| quad_perm:[0,1,2,3] row_mask:0x0 bank_mask:0x0 ; encoding: [0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x80,0x00]
+0xfa,0x04,0x0a,0x2c,0x01,0xe4,0x80,0x00
+
+# GFX950: v_dot2c_f32_bf16_e64 v5, v1, src_scc ; encoding: [0x05,0x00,0x16,0xd1,0x01,0xfb,0x01,0x00]
+0x05,0x00,0x16,0xd1,0x01,0xfb,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_e64 v5, v255, src_execz ; encoding: [0x05,0x00,0x16,0xd1,0xff,0xf9,0x01,0x00]
+0x05,0x00,0x16,0xd1,0xff,0xf9,0x01,0x00
+
+# GFX950: v_dot2c_f32_bf16_e64 v5, s101, s101 ; encoding: [0x05,0x00,0x16,0xd1,0x65,0xca,0x00,0x00]
+0x05,0x00,0x16,0xd1,0x65,0xca,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_e64 v5, -1, flat_scratch_lo ; encoding: [0x05,0x00,0x16,0xd1,0xc1,0xcc,0x00,0x00]
+0x05,0x00,0x16,0xd1,0xc1,0xcc,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_e64 v5, 0.5, -|flat_scratch_hi| ; encoding: [0x05,0x02,0x16,0xd1,0xf0,0xce,0x00,0x40]
+0x05,0x02,0x16,0xd1,0xf0,0xce,0x00,0x40
+
+# GFX950: v_dot2c_f32_bf16_e64 v5, src_execz, 0.5 mul:4 ; encoding: [0x05,0x00,0x16,0xd1,0xfc,0xe0,0x01,0x10]
+0x05,0x00,0x16,0xd1,0xfc,0xe0,0x01,0x10
+
+# GFX950: v_dot2c_f32_bf16_e64 v255, -|src_scc|, -1 clamp div:2 ; encoding: [0xff,0x81,0x16,0xd1,0xfd,0x82,0x01,0x38]
+0xff,0x81,0x16,0xd1,0xfd,0x82,0x01,0x38
+
+# GFX950: v_dot2c_f32_bf16_e32 v5, 10, v2 ; encoding: [0x8a,0x04,0x0a,0x2c]
+0x8a,0x04,0x0a,0x2c
+
+# GFX950: v_dot2c_f32_bf16_e32 v5, 0x64, v2 ; encoding: [0xff,0x04,0x0a,0x2c,0x64,0x00,0x00,0x00]
+0xff,0x04,0x0a,0x2c,0x64,0x00,0x00,0x00
+
+# GFX950: v_dot2c_f32_bf16_e32 v5, 0x4122, v2 ; encoding: [0xff,0x04,0x0a,0x2c,0x22,0x41,0x00,0x00]
+0xff,0x04,0x0a,0x2c,0x22,0x41,0x00,0x00
More information about the cfe-commits
mailing list