Skip to content

Commit 02ca83a

Browse files
shiltianshiltian
committed
[AMDGPU] Introduce a new generic target gfx9-4-generic 
1 parent aa79412 commit 02ca83a

35 files changed

+2257
-10
lines changed

‎clang/include/clang/Basic/Cuda.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -103,6 +103,7 @@ enum class OffloadArch {
103103
GFX909,
104104
GFX90a,
105105
GFX90c,
106+
GFX9_4_GENERIC,
106107
GFX940,
107108
GFX941,
108109
GFX942,

‎clang/lib/Basic/Cuda.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -121,6 +121,7 @@ static const OffloadArchToStringMap arch_names[] = {
121121
GFX(909), // gfx909
122122
GFX(90a), // gfx90a
123123
GFX(90c), // gfx90c
124+
{OffloadArch::GFX9_4_GENERIC, "gfx9-4-generic", "compute_amdgcn"},
124125
GFX(940), // gfx940
125126
GFX(941), // gfx941
126127
GFX(942), // gfx942

‎clang/lib/Basic/Targets/NVPTX.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -205,6 +205,7 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
205205
case OffloadArch::GFX909:
206206
case OffloadArch::GFX90a:
207207
case OffloadArch::GFX90c:
208+
case OffloadArch::GFX9_4_GENERIC:
208209
case OffloadArch::GFX940:
209210
case OffloadArch::GFX941:
210211
case OffloadArch::GFX942:

‎clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2301,6 +2301,7 @@ void CGOpenMPRuntimeGPU::processRequiresDirective(const OMPRequiresDecl *D) {
23012301
case OffloadArch::GFX909:
23022302
case OffloadArch::GFX90a:
23032303
case OffloadArch::GFX90c:
2304+
case OffloadArch::GFX9_4_GENERIC:
23042305
case OffloadArch::GFX940:
23052306
case OffloadArch::GFX941:
23062307
case OffloadArch::GFX942:

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,8 @@
5656

5757
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1103 -target-feature +wavefrontsize64 -emit-llvm -o - %s | FileCheck --check-prefix=GFX1103-W64 %s
5858

59+
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx9-4-generic -emit-llvm -o - %s | FileCheck --check-prefix=GFX9_4_Generic %s
60+
5961
// NOCPU-NOT: "target-features"
6062
// NOCPU-WAVE32: "target-features"="+wavefrontsize32"
6163
// NOCPU-WAVE64: "target-features"="+wavefrontsize64"
@@ -85,6 +87,7 @@
8587
// GFX940: "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"
8688
// 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"
8789
// 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"
90+
// 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,+fp8-conversion-insts,+fp8-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+mai-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64"
8891
// GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize32"
8992
// 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"
9093
// 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/Driver/amdgpu-macros.cl

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -133,6 +133,7 @@
133133
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1201 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx1201 -DFAMILY=GFX12
134134

135135
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx9-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx9_generic -DFAMILY=GFX9
136+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx9-4-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=64 -DCPU=gfx9_generic -DFAMILY=GFX9
136137
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx10-1-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx10_1_generic -DFAMILY=GFX10
137138
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx10-3-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx10_3_generic -DFAMILY=GFX10
138139
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx11-generic %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,FAST_FMAF %s -DWAVEFRONT_SIZE=32 -DCPU=gfx11_generic -DFAMILY=GFX11

‎clang/test/Driver/amdgpu-mcpu.cl

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -118,6 +118,7 @@
118118
// RUN: %clang -### -target amdgcn -mcpu=gfx1201 %s 2>&1 | FileCheck --check-prefix=GFX1201 %s
119119

120120
// RUN: %clang -### -target amdgcn -mcpu=gfx9-generic %s 2>&1 | FileCheck --check-prefix=GFX9_GENERIC %s
121+
// RUN: %clang -### -target amdgcn -mcpu=gfx9-4-generic %s 2>&1 | FileCheck --check-prefix=GFX9_4_GENERIC %s
121122
// RUN: %clang -### -target amdgcn -mcpu=gfx10-1-generic %s 2>&1 | FileCheck --check-prefix=GFX10_1_GENERIC %s
122123
// RUN: %clang -### -target amdgcn -mcpu=gfx10-3-generic %s 2>&1 | FileCheck --check-prefix=GFX10_3_GENERIC %s
123124
// RUN: %clang -### -target amdgcn -mcpu=gfx11-generic %s 2>&1 | FileCheck --check-prefix=GFX11_GENERIC %s
@@ -172,6 +173,7 @@
172173
// GFX1201: "-target-cpu" "gfx1201"
173174

174175
// GFX9_GENERIC: "-target-cpu" "gfx9-generic"
176+
// GFX9_4_GENERIC: "-target-cpu" "gfx9-4-generic"
175177
// GFX10_1_GENERIC: "-target-cpu" "gfx10-1-generic"
176178
// GFX10_3_GENERIC: "-target-cpu" "gfx10-3-generic"
177179
// GFX11_GENERIC: "-target-cpu" "gfx11-generic"

‎clang/test/Misc/target-invalid-cpu-note/amdgcn.c

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -70,6 +70,7 @@
7070
// CHECK-SAME: {{^}}, gfx1200
7171
// CHECK-SAME: {{^}}, gfx1201
7272
// CHECK-SAME: {{^}}, gfx9-generic
73+
// CHECK-SAME: {{^}}, gfx9-4-generic
7374
// CHECK-SAME: {{^}}, gfx10-1-generic
7475
// CHECK-SAME: {{^}}, gfx10-3-generic
7576
// CHECK-SAME: {{^}}, gfx11-generic

‎clang/test/Misc/target-invalid-cpu-note/nvptx.c

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,7 @@
5050
// CHECK-SAME: {{^}}, gfx909
5151
// CHECK-SAME: {{^}}, gfx90a
5252
// CHECK-SAME: {{^}}, gfx90c
53+
// CHECK-SAME: {{^}}, gfx9-4-generic
5354
// CHECK-SAME: {{^}}, gfx940
5455
// CHECK-SAME: {{^}}, gfx941
5556
// CHECK-SAME: {{^}}, gfx942

‎llvm/docs/AMDGPUUsage.rst

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -576,6 +576,17 @@ Generic processor code objects are versioned. See :ref:`amdgpu-generic-processor
576576
- ``v_dot2_f32_f16``
577577

578578

579+
``gfx9-4-generic`` ``amdgcn`` - ``gfx940`` - xnack - Absolute flat - FP8 related instructions are not available.
580+
- ``gfx941`` scratch - The following instructions are not available:
581+
- ``gfx942``
582+
- ``v_mfma_f32_16x16x8_xf32``
583+
- ``v_mfma_f32_32x32x4xf32``
584+
- ``v_cvt_f32_fp8``
585+
- ``v_cvt_f32_bf8``
586+
- ``v_cvt_pk_f32_fp8``
587+
- ``v_cvt_pk_f32_bf8``
588+
589+
579590
``gfx10-1-generic`` ``amdgcn`` - ``gfx1010`` - xnack - Absolute flat - The following instructions are
580591
- ``gfx1011`` - wavefrontsize64 scratch not available on ``gfx1011``
581592
- ``gfx1012`` - cumode and ``gfx1012``

‎llvm/include/llvm/BinaryFormat/ELF.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -822,11 +822,12 @@ enum : unsigned {
822822
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X57 = 0x057,
823823
EF_AMDGPU_MACH_AMDGCN_GFX1153 = 0x058,
824824
EF_AMDGPU_MACH_AMDGCN_GFX12_GENERIC = 0x059,
825+
EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC = 0x05f,
825826
// clang-format on
826827

827828
// First/last AMDGCN-based processors.
828829
EF_AMDGPU_MACH_AMDGCN_FIRST = EF_AMDGPU_MACH_AMDGCN_GFX600,
829-
EF_AMDGPU_MACH_AMDGCN_LAST = EF_AMDGPU_MACH_AMDGCN_GFX12_GENERIC,
830+
EF_AMDGPU_MACH_AMDGCN_LAST = EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC,
830831

831832
// Indicates if the "xnack" target feature is enabled for all code contained
832833
// in the object.

‎llvm/include/llvm/TargetParser/TargetParser.h

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -119,9 +119,10 @@ enum GPUKind : uint32_t {
119119
GK_GFX10_3_GENERIC = 194,
120120
GK_GFX11_GENERIC = 195,
121121
GK_GFX12_GENERIC = 196,
122+
GK_GFX9_4_GENERIC = 197,
122123

123124
GK_AMDGCN_GENERIC_FIRST = GK_GFX9_GENERIC,
124-
GK_AMDGCN_GENERIC_LAST = GK_GFX12_GENERIC,
125+
GK_AMDGCN_GENERIC_LAST = GK_GFX9_4_GENERIC,
125126
};
126127

127128
/// Instruction set architecture version.

‎llvm/lib/Object/ELFObjectFile.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -602,6 +602,8 @@ StringRef ELFObjectFileBase::getAMDGPUCPUName() const {
602602
// Generic AMDGCN targets
603603
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC:
604604
return "gfx9-generic";
605+
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC:
606+
return "gfx9-4-generic";
605607
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX10_1_GENERIC:
606608
return "gfx10-1-generic";
607609
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX10_3_GENERIC:

‎llvm/lib/ObjectYAML/ELFYAML.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -631,6 +631,7 @@ void ScalarBitSetTraits<ELFYAML::ELF_EF>::bitset(IO &IO,
631631
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX1200, EF_AMDGPU_MACH);
632632
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX1201, EF_AMDGPU_MACH);
633633
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC, EF_AMDGPU_MACH);
634+
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC, EF_AMDGPU_MACH);
634635
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX10_1_GENERIC, EF_AMDGPU_MACH);
635636
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX10_3_GENERIC, EF_AMDGPU_MACH);
636637
BCaseMask(EF_AMDGPU_MACH_AMDGCN_GFX11_GENERIC, EF_AMDGPU_MACH);

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

Lines changed: 33 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -719,6 +719,12 @@ def FeatureAtomicFaddRtnInsts : SubtargetFeature<"atomic-fadd-rtn-insts",
719719
[FeatureFlatGlobalInsts]
720720
>;
721721

722+
def FeatureXF32Insts : SubtargetFeature<"xf32-insts",
723+
"HasXF32Insts",
724+
"true",
725+
"Has v_mfma_f32_16x16x8xf32 and v_mfma_f32_32x32x4xf32 instructions"
726+
>;
727+
722728
def FeatureAtomicFMinFMaxF32GlobalInsts : SubtargetFeature<"atomic-fmin-fmax-global-f32",
723729
"HasAtomicFMinFMaxF32GlobalInsts",
724730
"true",
@@ -1444,9 +1450,6 @@ def FeatureISAVersion9_4_Common : FeatureSet<
14441450
FeatureDPALU_DPP,
14451451
FeaturePackedFP32Ops,
14461452
FeatureMAIInsts,
1447-
FeatureFP8Insts,
1448-
FeatureFP8ConversionInsts,
1449-
FeatureCvtFP8VOP1Bug,
14501453
FeaturePkFmacF16Inst,
14511454
FeatureAtomicFaddRtnInsts,
14521455
FeatureAtomicFaddNoRtnInsts,
@@ -1468,15 +1471,36 @@ def FeatureISAVersion9_4_Common : FeatureSet<
14681471

14691472
def FeatureISAVersion9_4_0 : FeatureSet<
14701473
!listconcat(FeatureISAVersion9_4_Common.Features,
1471-
[FeatureForceStoreSC0SC1])>;
1474+
[
1475+
FeatureForceStoreSC0SC1,
1476+
FeatureFP8Insts,
1477+
FeatureFP8ConversionInsts,
1478+
FeatureCvtFP8VOP1Bug,
1479+
FeatureXF32Insts
1480+
])>;
14721481

14731482
def FeatureISAVersion9_4_1 : FeatureSet<
14741483
!listconcat(FeatureISAVersion9_4_Common.Features,
1475-
[FeatureForceStoreSC0SC1])>;
1484+
[
1485+
FeatureForceStoreSC0SC1,
1486+
FeatureFP8Insts,
1487+
FeatureFP8ConversionInsts,
1488+
FeatureCvtFP8VOP1Bug,
1489+
FeatureXF32Insts
1490+
])>;
14761491

14771492
def FeatureISAVersion9_4_2 : FeatureSet<
14781493
!listconcat(FeatureISAVersion9_4_Common.Features,
1479-
[])>;
1494+
[
1495+
FeatureFP8Insts,
1496+
FeatureFP8ConversionInsts,
1497+
FeatureCvtFP8VOP1Bug,
1498+
FeatureXF32Insts
1499+
])>;
1500+
1501+
def FeatureISAVersion9_4_Generic : FeatureSet<
1502+
!listconcat(FeatureISAVersion9_4_Common.Features,
1503+
[FeatureRequiresCOV6])>;
14801504

14811505
def FeatureISAVersion10_Common : FeatureSet<
14821506
[FeatureGFX10,
@@ -2021,6 +2045,9 @@ def HasRestrictedSOffset : Predicate<"Subtarget->hasRestrictedSOffset()">,
20212045
def HasUnrestrictedSOffset : Predicate<"!Subtarget->hasRestrictedSOffset()">,
20222046
AssemblerPredicate<(all_of (not FeatureHasRestrictedSOffset))>;
20232047

2048+
def HasXF32Insts : Predicate<"Subtarget->hasXF32Insts()">,
2049+
AssemblerPredicate<(all_of FeatureXF32Insts)>;
2050+
20242051
def D16PreservesUnusedBits :
20252052
Predicate<"Subtarget->d16PreservesUnusedBits()">,
20262053
AssemblerPredicate<(all_of FeatureGFX9Insts, (not FeatureSRAMECC))>;

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -209,6 +209,11 @@ def : ProcessorModel<"gfx9-generic", SIQuarterSpeedModel,
209209
FeatureISAVersion9_Generic.Features
210210
>;
211211

212+
// [gfx940, gfx941, gfx942]
213+
def : ProcessorModel<"gfx9-4-generic", SIQuarterSpeedModel,
214+
FeatureISAVersion9_4_Generic.Features
215+
>;
216+
212217
//===----------------------------------------------------------------------===//
213218
// GCN GFX10.
214219
//===----------------------------------------------------------------------===//

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -179,6 +179,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
179179
bool HasDefaultComponentZero = false;
180180
bool HasAgentScopeFineGrainedRemoteMemoryAtomics = false;
181181
bool HasDefaultComponentBroadcast = false;
182+
bool HasXF32Insts = false;
182183
/// The maximum number of instructions that may be placed within an S_CLAUSE,
183184
/// which is one greater than the maximum argument to S_CLAUSE. A value of 0
184185
/// indicates a lack of S_CLAUSE support.
@@ -1302,6 +1303,8 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
13021303
return getGeneration() == GFX12;
13031304
}
13041305

1306+
bool hasXF32Insts() const { return HasXF32Insts; }
1307+
13051308
/// \returns The maximum number of instructions that can be enclosed in an
13061309
/// S_CLAUSE on the given subtarget, or 0 for targets that do not support that
13071310
/// instruction.

‎llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -119,6 +119,7 @@ StringRef AMDGPUTargetStreamer::getArchNameFromElfMach(unsigned ElfMach) {
119119
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX1200: AK = GK_GFX1200; break;
120120
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX1201: AK = GK_GFX1201; break;
121121
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC: AK = GK_GFX9_GENERIC; break;
122+
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC: AK = GK_GFX9_4_GENERIC; break;
122123
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX10_1_GENERIC: AK = GK_GFX10_1_GENERIC; break;
123124
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX10_3_GENERIC: AK = GK_GFX10_3_GENERIC; break;
124125
case ELF::EF_AMDGPU_MACH_AMDGCN_GFX11_GENERIC: AK = GK_GFX11_GENERIC; break;
@@ -204,6 +205,7 @@ unsigned AMDGPUTargetStreamer::getElfMach(StringRef GPU) {
204205
case GK_GFX1200: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX1200;
205206
case GK_GFX1201: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX1201;
206207
case GK_GFX9_GENERIC: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC;
208+
case GK_GFX9_4_GENERIC: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC;
207209
case GK_GFX10_1_GENERIC: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX10_1_GENERIC;
208210
case GK_GFX10_3_GENERIC: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX10_3_GENERIC;
209211
case GK_GFX11_GENERIC: return ELF::EF_AMDGPU_MACH_AMDGCN_GFX11_GENERIC;
@@ -821,6 +823,9 @@ unsigned AMDGPUTargetELFStreamer::getEFlagsV6() {
821823
case AMDGPU::GK_GFX9_GENERIC:
822824
Version = GenericVersion::GFX9;
823825
break;
826+
case AMDGPU::GK_GFX9_4_GENERIC:
827+
Version = GenericVersion::GFX9_4;
828+
break;
824829
case AMDGPU::GK_GFX10_1_GENERIC:
825830
Version = GenericVersion::GFX10_1;
826831
break;

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,7 @@ struct IsaVersion;
4646
/// within a generic family.
4747
namespace GenericVersion {
4848
static constexpr unsigned GFX9 = 1;
49+
static constexpr unsigned GFX9_4 = 1;
4950
static constexpr unsigned GFX10_1 = 1;
5051
static constexpr unsigned GFX10_3 = 1;
5152
static constexpr unsigned GFX11 = 1;

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

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -757,8 +757,6 @@ let Predicates = [isGFX90APlus] in {
757757
let SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1 in {
758758
defm V_MFMA_I32_32X32X16I8 : MAIInst<"v_mfma_i32_32x32x16i8", "I32_I64_X32", int_amdgcn_mfma_i32_32x32x16_i8>;
759759
defm V_MFMA_I32_16X16X32I8 : MAIInst<"v_mfma_i32_16x16x32i8", "I32_I64_X16", int_amdgcn_mfma_i32_16x16x32_i8>;
760-
defm V_MFMA_F32_16X16X8XF32 : MAIInst<"v_mfma_f32_16x16x8xf32", "F32_V2F32_X16", int_amdgcn_mfma_f32_16x16x8_xf32>;
761-
defm V_MFMA_F32_32X32X4XF32 : MAIInst<"v_mfma_f32_32x32x4xf32", "F32_V2F32_X32", int_amdgcn_mfma_f32_32x32x4_xf32>;
762760
defm V_MFMA_F32_16X16X32_BF8_BF8 : MAIInst<"v_mfma_f32_16x16x32_bf8_bf8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_bf8_bf8>;
763761
defm V_MFMA_F32_16X16X32_BF8_FP8 : MAIInst<"v_mfma_f32_16x16x32_bf8_fp8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_bf8_fp8>;
764762
defm V_MFMA_F32_16X16X32_FP8_BF8 : MAIInst<"v_mfma_f32_16x16x32_fp8_bf8", "F32_I64_X32", int_amdgcn_mfma_f32_16x16x32_fp8_bf8>;
@@ -769,6 +767,11 @@ let SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1 in {
769767
defm V_MFMA_F32_32X32X16_FP8_FP8 : MAIInst<"v_mfma_f32_32x32x16_fp8_fp8", "F32_I64_X16", int_amdgcn_mfma_f32_32x32x16_fp8_fp8>;
770768
} // End SubtargetPredicate = isGFX940Plus, is_gfx940_xdl = 1
771769

770+
let SubtargetPredicate = HasXF32Insts, is_gfx940_xdl = 1 in {
771+
defm V_MFMA_F32_16X16X8XF32 : MAIInst<"v_mfma_f32_16x16x8xf32", "F32_V2F32_X16", int_amdgcn_mfma_f32_16x16x8_xf32>;
772+
defm V_MFMA_F32_32X32X4XF32 : MAIInst<"v_mfma_f32_32x32x4xf32", "F32_V2F32_X32", int_amdgcn_mfma_f32_32x32x4_xf32>;
773+
} // End SubtargetPredicate = HasXF32Insts, is_gfx940_xdl = 1
774+
772775
multiclass SMFMACInst<string OpName, string P, SDPatternOperator node> {
773776
let Constraints = "$vdst = $src2", DisableEncoding = "$src2",
774777
isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1, is_gfx940_xdl = 1 in {
@@ -1757,8 +1760,10 @@ defm V_MFMA_F64_4X4X4F64 : VOP3P_Real_MFMA_gfx90a <0x6f>;
17571760

17581761
defm V_MFMA_I32_32X32X16I8 : VOP3P_Real_MFMA_gfx940 <0x56, "v_mfma_i32_32x32x16_i8">;
17591762
defm V_MFMA_I32_16X16X32I8 : VOP3P_Real_MFMA_gfx940 <0x57, "v_mfma_i32_16x16x32_i8">;
1763+
let SubtargetPredicate = HasXF32Insts in {
17601764
defm V_MFMA_F32_16X16X8XF32 : VOP3P_Real_MFMA_gfx940 <0x3e, "v_mfma_f32_16x16x8_xf32">;
17611765
defm V_MFMA_F32_32X32X4XF32 : VOP3P_Real_MFMA_gfx940 <0x3f, "v_mfma_f32_32x32x4_xf32">;
1766+
} // End SubtargetPredicate = HasXF32Insts
17621767
defm V_MFMA_F32_16X16X32_BF8_BF8 : VOP3P_Real_MFMA_gfx940 <0x70>;
17631768
defm V_MFMA_F32_16X16X32_BF8_FP8 : VOP3P_Real_MFMA_gfx940 <0x71>;
17641769
defm V_MFMA_F32_16X16X32_FP8_BF8 : VOP3P_Real_MFMA_gfx940 <0x72>;

‎llvm/lib/TargetParser/TargetParser.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -130,6 +130,7 @@ constexpr GPUInfo AMDGCNGPUs[] = {
130130
{{"gfx1201"}, {"gfx1201"}, GK_GFX1201, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32|FEATURE_WGP},
131131

132132
{{"gfx9-generic"}, {"gfx9-generic"}, GK_GFX9_GENERIC, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_XNACK},
133+
{{"gfx9-4-generic"}, {"gfx9-4-generic"}, GK_GFX9_4_GENERIC, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_XNACK|FEATURE_SRAMECC},
133134
{{"gfx10-1-generic"}, {"gfx10-1-generic"}, GK_GFX10_1_GENERIC, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32|FEATURE_XNACK|FEATURE_WGP},
134135
{{"gfx10-3-generic"}, {"gfx10-3-generic"}, GK_GFX10_3_GENERIC, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32|FEATURE_WGP},
135136
{{"gfx11-generic"}, {"gfx11-generic"}, GK_GFX11_GENERIC, FEATURE_FAST_FMA_F32|FEATURE_FAST_DENORMAL_F32|FEATURE_WAVE32|FEATURE_WGP},
@@ -156,6 +157,8 @@ StringRef llvm::AMDGPU::getArchFamilyNameAMDGCN(GPUKind AK) {
156157
switch (AK) {
157158
case AMDGPU::GK_GFX9_GENERIC:
158159
return "gfx9";
160+
case AMDGPU::GK_GFX9_4_GENERIC:
161+
return "gfx9";
159162
case AMDGPU::GK_GFX10_1_GENERIC:
160163
case AMDGPU::GK_GFX10_3_GENERIC:
161164
return "gfx10";
@@ -296,6 +299,7 @@ AMDGPU::IsaVersion AMDGPU::getIsaVersion(StringRef GPU) {
296299
// TODO: Split up this API depending on its caller so
297300
// generic target handling is more obvious and less risky.
298301
case GK_GFX9_GENERIC: return {9, 0, 0};
302+
case GK_GFX9_4_GENERIC: return {9, 4, 0};
299303
case GK_GFX10_1_GENERIC: return {10, 1, 0};
300304
case GK_GFX10_3_GENERIC: return {10, 3, 0};
301305
case GK_GFX11_GENERIC: return {11, 0, 3};
@@ -466,6 +470,7 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const Triple &T,
466470
case GK_GFX942:
467471
case GK_GFX941:
468472
case GK_GFX940:
473+
case GK_GFX9_4_GENERIC:
469474
Features["gfx940-insts"] = true;
470475
Features["fp8-insts"] = true;
471476
Features["fp8-conversion-insts"] = true;

0 commit comments

Comments
 (0)