Skip to content

Commit d2da3c2

Browse files
arsenmarsenm
committed
AMDGPU: Add Vega12 and Vega20
Changes by Matt Arsenault Konstantin Zhuravlyov llvm-svn: 331216
1 parent 0084adc commit d2da3c2

File tree

8 files changed

+113
-2
lines changed

8 files changed

+113
-2
lines changed

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -120,6 +120,18 @@ TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
120120

121121
TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
122122

123+
//===----------------------------------------------------------------------===//
124+
// Deep learning builtins.
125+
//===----------------------------------------------------------------------===//
126+
127+
TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hf", "nc", "dl-insts")
128+
TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSi", "nc", "dl-insts")
129+
TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUi", "nc", "dl-insts")
130+
TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSi", "nc", "dl-insts")
131+
TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUi", "nc", "dl-insts")
132+
TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSi", "nc", "dl-insts")
133+
TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUi", "nc", "dl-insts")
134+
123135
//===----------------------------------------------------------------------===//
124136
// Special builtins.
125137
//===----------------------------------------------------------------------===//

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -133,6 +133,10 @@ bool AMDGPUTargetInfo::initFeatureMap(
133133
CPU = "gfx600";
134134

135135
switch (parseAMDGCNName(CPU).Kind) {
136+
case GK_GFX906:
137+
Features["dl-insts"] = true;
138+
LLVM_FALLTHROUGH;
139+
case GK_GFX904:
136140
case GK_GFX902:
137141
case GK_GFX900:
138142
Features["gfx9-insts"] = true;

‎clang/lib/Basic/Targets/AMDGPU.h‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -78,9 +78,11 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
7878
GK_GFX810,
7979
GK_GFX900,
8080
GK_GFX902,
81+
GK_GFX904,
82+
GK_GFX906,
8183

8284
GK_AMDGCN_FIRST = GK_GFX600,
83-
GK_AMDGCN_LAST = GK_GFX902,
85+
GK_AMDGCN_LAST = GK_GFX906,
8486
};
8587

8688
struct GPUInfo {
@@ -127,7 +129,7 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
127129
{{"cayman"}, {"cayman"}, GK_CAYMAN, true, false, false, false, false},
128130
{{"turks"}, {"turks"}, GK_TURKS, false, false, false, false, false},
129131
};
130-
static constexpr GPUInfo AMDGCNGPUs[30] = {
132+
static constexpr GPUInfo AMDGCNGPUs[32] = {
131133
// Name Canonical Kind Has Has Has Has Has
132134
// Name FMAF Fast LDEXPF FP64 Fast
133135
// FMAF FMA
@@ -161,6 +163,8 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
161163
{{"stoney"}, {"gfx810"}, GK_GFX810, true, false, true, true, true},
162164
{{"gfx900"}, {"gfx900"}, GK_GFX900, true, true, true, true, true},
163165
{{"gfx902"}, {"gfx902"}, GK_GFX900, true, true, true, true, true},
166+
{{"gfx904"}, {"gfx904"}, GK_GFX904, true, true, true, true, true},
167+
{{"gfx906"}, {"gfx906"}, GK_GFX906, true, true, true, true, true},
164168
};
165169

166170
static GPUInfo parseR600Name(StringRef Name);

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// REQUIRES: amdgpu-registered-target
2+
3+
// Check that appropriate features are defined for every supported AMDGPU
4+
// "-target" and "-mcpu" options.
5+
6+
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
7+
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
8+
9+
// GFX904: "target-features"="+16-bit-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
10+
// GFX906: "target-features"="+16-bit-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
11+
12+
kernel void test() {}

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

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// REQUIRES: amdgpu-registered-target
2+
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -S -emit-llvm -o - %s
4+
5+
typedef unsigned int uint;
6+
typedef half __attribute__((ext_vector_type(2))) half2;
7+
typedef short __attribute__((ext_vector_type(2))) short2;
8+
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
9+
10+
kernel void builtins_amdgcn_dl_insts_err(
11+
global float *fOut, global int *siOut, global uint *uiOut,
12+
half2 v2hA, half2 v2hB, float fC,
13+
short2 v2ssA, short2 v2ssB, int siA, int siB, int siC,
14+
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC) {
15+
fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC); // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dl-insts}}
16+
17+
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dl-insts}}
18+
uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC); // expected-error {{'__builtin_amdgcn_udot2' needs target feature dl-insts}}
19+
20+
siOut[1] = __builtin_amdgcn_sdot4(siA, siB, siC); // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dl-insts}}
21+
uiOut[1] = __builtin_amdgcn_udot4(uiA, uiB, uiC); // expected-error {{'__builtin_amdgcn_udot4' needs target feature dl-insts}}
22+
23+
siOut[2] = __builtin_amdgcn_sdot8(siA, siB, siC); // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dl-insts}}
24+
uiOut[2] = __builtin_amdgcn_udot8(uiA, uiB, uiC); // expected-error {{'__builtin_amdgcn_udot8' needs target feature dl-insts}}
25+
}

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

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// REQUIRES: amdgpu-registered-target
2+
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
4+
5+
typedef unsigned int uint;
6+
typedef half __attribute__((ext_vector_type(2))) half2;
7+
typedef short __attribute__((ext_vector_type(2))) short2;
8+
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
9+
10+
// CHECK-LABEL: @builtins_amdgcn_dl_insts
11+
// CHECK: call float @llvm.amdgcn.fdot2
12+
13+
// CHECK: call i32 @llvm.amdgcn.sdot2
14+
// CHECK: call i32 @llvm.amdgcn.udot2
15+
16+
// CHECK: call i32 @llvm.amdgcn.sdot4
17+
// CHECK: call i32 @llvm.amdgcn.udot4
18+
19+
// CHECK: call i32 @llvm.amdgcn.sdot8
20+
// CHECK: call i32 @llvm.amdgcn.udot8
21+
kernel void builtins_amdgcn_dl_insts(
22+
global float *fOut, global int *siOut, global uint *uiOut,
23+
half2 v2hA, half2 v2hB, float fC,
24+
short2 v2ssA, short2 v2ssB, int siA, int siB, int siC,
25+
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC) {
26+
fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC);
27+
28+
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC);
29+
uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC);
30+
31+
siOut[1] = __builtin_amdgcn_sdot4(siA, siB, siC);
32+
uiOut[1] = __builtin_amdgcn_udot4(uiA, uiB, uiC);
33+
34+
siOut[2] = __builtin_amdgcn_sdot8(siA, siB, siC);
35+
uiOut[2] = __builtin_amdgcn_udot8(uiA, uiB, uiC);
36+
}

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

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -173,6 +173,8 @@
173173
// RUN: %clang -E -dM -target amdgcn -mcpu=stoney %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX810 %s
174174
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX900 %s
175175
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX902 %s
176+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX904 %s
177+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX906 %s
176178

177179
// GFX600-DAG: #define FP_FAST_FMA 1
178180
// GFX601-DAG: #define FP_FAST_FMA 1
@@ -187,6 +189,8 @@
187189
// GFX810-DAG: #define FP_FAST_FMA 1
188190
// GFX900-DAG: #define FP_FAST_FMA 1
189191
// GFX902-DAG: #define FP_FAST_FMA 1
192+
// GFX904-DAG: #define FP_FAST_FMA 1
193+
// GFX906-DAG: #define FP_FAST_FMA 1
190194

191195
// GFX600-DAG: #define FP_FAST_FMAF 1
192196
// GFX601-NOT: #define FP_FAST_FMAF 1
@@ -201,6 +205,8 @@
201205
// GFX810-NOT: #define FP_FAST_FMAF 1
202206
// GFX900-DAG: #define FP_FAST_FMAF 1
203207
// GFX902-DAG: #define FP_FAST_FMAF 1
208+
// GFX904-DAG: #define FP_FAST_FMAF 1
209+
// GFX906-DAG: #define FP_FAST_FMAF 1
204210

205211
// ARCH-GCN-DAG: #define __AMDGCN__ 1
206212
// ARCH-GCN-DAG: #define __AMDGPU__ 1
@@ -219,6 +225,8 @@
219225
// GFX810-DAG: #define __HAS_FMAF__ 1
220226
// GFX900-DAG: #define __HAS_FMAF__ 1
221227
// GFX902-DAG: #define __HAS_FMAF__ 1
228+
// GFX904-DAG: #define __HAS_FMAF__ 1
229+
// GFX906-DAG: #define __HAS_FMAF__ 1
222230

223231
// GFX600-DAG: #define __HAS_FP64__ 1
224232
// GFX601-DAG: #define __HAS_FP64__ 1
@@ -233,6 +241,8 @@
233241
// GFX810-DAG: #define __HAS_FP64__ 1
234242
// GFX900-DAG: #define __HAS_FP64__ 1
235243
// GFX902-DAG: #define __HAS_FP64__ 1
244+
// GFX904-DAG: #define __HAS_FP64__ 1
245+
// GFX906-DAG: #define __HAS_FP64__ 1
236246

237247
// GFX600-DAG: #define __HAS_LDEXPF__ 1
238248
// GFX601-DAG: #define __HAS_LDEXPF__ 1
@@ -247,6 +257,8 @@
247257
// GFX810-DAG: #define __HAS_LDEXPF__ 1
248258
// GFX900-DAG: #define __HAS_LDEXPF__ 1
249259
// GFX902-DAG: #define __HAS_LDEXPF__ 1
260+
// GFX904-DAG: #define __HAS_LDEXPF__ 1
261+
// GFX906-DAG: #define __HAS_LDEXPF__ 1
250262

251263
// GFX600-DAG: #define __gfx600__ 1
252264
// GFX601-DAG: #define __gfx601__ 1
@@ -261,3 +273,5 @@
261273
// GFX810-DAG: #define __gfx810__ 1
262274
// GFX900-DAG: #define __gfx900__ 1
263275
// GFX902-DAG: #define __gfx902__ 1
276+
// GFX904-DAG: #define __gfx904__ 1
277+
// GFX906-DAG: #define __gfx906__ 1

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,8 @@
8282
// RUN: %clang -### -target amdgcn -mcpu=stoney %s 2>&1 | FileCheck --check-prefix=STONEY %s
8383
// RUN: %clang -### -target amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck --check-prefix=GFX900 %s
8484
// RUN: %clang -### -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefix=GFX902 %s
85+
// RUN: %clang -### -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefix=GFX904 %s
86+
// RUN: %clang -### -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefix=GFX906 %s
8587

8688
// GFX600: "-target-cpu" "gfx600"
8789
// TAHITI: "-target-cpu" "tahiti"
@@ -113,3 +115,5 @@
113115
// STONEY: "-target-cpu" "stoney"
114116
// GFX900: "-target-cpu" "gfx900"
115117
// GFX902: "-target-cpu" "gfx902"
118+
// GFX904: "-target-cpu" "gfx904"
119+
// GFX906: "-target-cpu" "gfx906"

0 commit comments

Comments
 (0)