Skip to content

Commit 02ca83a

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

File tree

35 files changed

+2257
-10
lines changed

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``

0 commit comments

Comments
 (0)