maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

AMDGPU: Add Vega12 and Vega20 

Changes by
  Matt Arsenault
  Konstantin Zhuravlyov

llvm-svn: 331216
This commit is contained in:
Matt Arsenault 2018-04-30 19:08:27 +00:00
parent 0084adc516
commit d2da3c20d7
8 changed files with 113 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -120,6 +120,18 @@ TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
//===----------------------------------------------------------------------===//
// Deep learning builtins.
//===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hf", "nc", "dl-insts")
TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSi", "nc", "dl-insts")
TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUi", "nc", "dl-insts")
TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSi", "nc", "dl-insts")
TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUi", "nc", "dl-insts")
TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSi", "nc", "dl-insts")
TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUi", "nc", "dl-insts")
//===----------------------------------------------------------------------===//
// Special builtins.
//===----------------------------------------------------------------------===//

4
clang/lib/Basic/Targets/AMDGPU.cpp
View File

@ -133,6 +133,10 @@ bool AMDGPUTargetInfo::initFeatureMap(
CPU = "gfx600";
switch (parseAMDGCNName(CPU).Kind) {
case GK_GFX906:
Features["dl-insts"] = true;
LLVM_FALLTHROUGH;
case GK_GFX904:
case GK_GFX902:
case GK_GFX900:
Features["gfx9-insts"] = true;

8
clang/lib/Basic/Targets/AMDGPU.h
View File

@ -78,9 +78,11 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
GK_GFX810,
GK_GFX900,
GK_GFX902,
GK_GFX904,
GK_GFX906,
GK_AMDGCN_FIRST = GK_GFX600,
GK_AMDGCN_LAST = GK_GFX902,
GK_AMDGCN_LAST = GK_GFX906,
};
struct GPUInfo {
@ -127,7 +129,7 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
{{"cayman"}, {"cayman"}, GK_CAYMAN, true, false, false, false, false},
{{"turks"}, {"turks"}, GK_TURKS, false, false, false, false, false},
};
static constexpr GPUInfo AMDGCNGPUs[30] = {
static constexpr GPUInfo AMDGCNGPUs[32] = {
// Name Canonical Kind Has Has Has Has Has
// Name FMAF Fast LDEXPF FP64 Fast
// FMAF FMA
@ -161,6 +163,8 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
{{"stoney"}, {"gfx810"}, GK_GFX810, true, false, true, true, true},
{{"gfx900"}, {"gfx900"}, GK_GFX900, true, true, true, true, true},
{{"gfx902"}, {"gfx902"}, GK_GFX900, true, true, true, true, true},
{{"gfx904"}, {"gfx904"}, GK_GFX904, true, true, true, true, true},
{{"gfx906"}, {"gfx906"}, GK_GFX906, true, true, true, true, true},
};
static GPUInfo parseR600Name(StringRef Name);

12
clang/test/CodeGenOpenCL/amdgpu-features.cl Normal file
View File

@ -0,0 +1,12 @@
// REQUIRES: amdgpu-registered-target
// Check that appropriate features are defined for every supported AMDGPU
// "-target" and "-mcpu" options.
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
// GFX904: "target-features"="+16-bit-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
// GFX906: "target-features"="+16-bit-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
kernel void test() {}

25
clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl Normal file
View File

@ -0,0 +1,25 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -S -emit-llvm -o - %s
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;
kernel void builtins_amdgcn_dl_insts_err(
global float *fOut, global int *siOut, global uint *uiOut,
half2 v2hA, half2 v2hB, float fC,
short2 v2ssA, short2 v2ssB, int siA, int siB, int siC,
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC) {
fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC); // expected-error {{'__builtin_amdgcn_fdot2' needs target feature dl-insts}}
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC); // expected-error {{'__builtin_amdgcn_sdot2' needs target feature dl-insts}}
uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC); // expected-error {{'__builtin_amdgcn_udot2' needs target feature dl-insts}}
siOut[1] = __builtin_amdgcn_sdot4(siA, siB, siC); // expected-error {{'__builtin_amdgcn_sdot4' needs target feature dl-insts}}
uiOut[1] = __builtin_amdgcn_udot4(uiA, uiB, uiC); // expected-error {{'__builtin_amdgcn_udot4' needs target feature dl-insts}}
siOut[2] = __builtin_amdgcn_sdot8(siA, siB, siC); // expected-error {{'__builtin_amdgcn_sdot8' needs target feature dl-insts}}
uiOut[2] = __builtin_amdgcn_udot8(uiA, uiB, uiC); // expected-error {{'__builtin_amdgcn_udot8' needs target feature dl-insts}}
}

36
clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl Normal file
View File

@ -0,0 +1,36 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
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;
// CHECK-LABEL: @builtins_amdgcn_dl_insts
// CHECK: call float @llvm.amdgcn.fdot2
// CHECK: call i32 @llvm.amdgcn.sdot2
// CHECK: call i32 @llvm.amdgcn.udot2
// CHECK: call i32 @llvm.amdgcn.sdot4
// CHECK: call i32 @llvm.amdgcn.udot4
// CHECK: call i32 @llvm.amdgcn.sdot8
// CHECK: call i32 @llvm.amdgcn.udot8
kernel void builtins_amdgcn_dl_insts(
global float *fOut, global int *siOut, global uint *uiOut,
half2 v2hA, half2 v2hB, float fC,
short2 v2ssA, short2 v2ssB, int siA, int siB, int siC,
ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC) {
fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC);
siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC);
uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC);
siOut[1] = __builtin_amdgcn_sdot4(siA, siB, siC);
uiOut[1] = __builtin_amdgcn_udot4(uiA, uiB, uiC);
siOut[2] = __builtin_amdgcn_sdot8(siA, siB, siC);
uiOut[2] = __builtin_amdgcn_udot8(uiA, uiB, uiC);
}

14
clang/test/Driver/amdgpu-macros.cl
View File

@ -173,6 +173,8 @@
// RUN: %clang -E -dM -target amdgcn -mcpu=stoney %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX810 %s
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX900 %s
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX902 %s
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX904 %s
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefixes=ARCH-GCN,GFX906 %s
// GFX600-DAG: #define FP_FAST_FMA 1
// GFX601-DAG: #define FP_FAST_FMA 1
@ -187,6 +189,8 @@
// GFX810-DAG: #define FP_FAST_FMA 1
// GFX900-DAG: #define FP_FAST_FMA 1
// GFX902-DAG: #define FP_FAST_FMA 1
// GFX904-DAG: #define FP_FAST_FMA 1
// GFX906-DAG: #define FP_FAST_FMA 1
// GFX600-DAG: #define FP_FAST_FMAF 1
// GFX601-NOT: #define FP_FAST_FMAF 1
@ -201,6 +205,8 @@
// GFX810-NOT: #define FP_FAST_FMAF 1
// GFX900-DAG: #define FP_FAST_FMAF 1
// GFX902-DAG: #define FP_FAST_FMAF 1
// GFX904-DAG: #define FP_FAST_FMAF 1
// GFX906-DAG: #define FP_FAST_FMAF 1
// ARCH-GCN-DAG: #define __AMDGCN__ 1
// ARCH-GCN-DAG: #define __AMDGPU__ 1
@ -219,6 +225,8 @@
// GFX810-DAG: #define __HAS_FMAF__ 1
// GFX900-DAG: #define __HAS_FMAF__ 1
// GFX902-DAG: #define __HAS_FMAF__ 1
// GFX904-DAG: #define __HAS_FMAF__ 1
// GFX906-DAG: #define __HAS_FMAF__ 1
// GFX600-DAG: #define __HAS_FP64__ 1
// GFX601-DAG: #define __HAS_FP64__ 1
@ -233,6 +241,8 @@
// GFX810-DAG: #define __HAS_FP64__ 1
// GFX900-DAG: #define __HAS_FP64__ 1
// GFX902-DAG: #define __HAS_FP64__ 1
// GFX904-DAG: #define __HAS_FP64__ 1
// GFX906-DAG: #define __HAS_FP64__ 1
// GFX600-DAG: #define __HAS_LDEXPF__ 1
// GFX601-DAG: #define __HAS_LDEXPF__ 1
@ -247,6 +257,8 @@
// GFX810-DAG: #define __HAS_LDEXPF__ 1
// GFX900-DAG: #define __HAS_LDEXPF__ 1
// GFX902-DAG: #define __HAS_LDEXPF__ 1
// GFX904-DAG: #define __HAS_LDEXPF__ 1
// GFX906-DAG: #define __HAS_LDEXPF__ 1
// GFX600-DAG: #define __gfx600__ 1
// GFX601-DAG: #define __gfx601__ 1
@ -261,3 +273,5 @@
// GFX810-DAG: #define __gfx810__ 1
// GFX900-DAG: #define __gfx900__ 1
// GFX902-DAG: #define __gfx902__ 1
// GFX904-DAG: #define __gfx904__ 1
// GFX906-DAG: #define __gfx906__ 1

4
clang/test/Driver/amdgpu-mcpu.cl
View File

@ -82,6 +82,8 @@
// RUN: %clang -### -target amdgcn -mcpu=stoney %s 2>&1 | FileCheck --check-prefix=STONEY %s
// RUN: %clang -### -target amdgcn -mcpu=gfx900 %s 2>&1 | FileCheck --check-prefix=GFX900 %s
// RUN: %clang -### -target amdgcn -mcpu=gfx902 %s 2>&1 | FileCheck --check-prefix=GFX902 %s
// RUN: %clang -### -target amdgcn -mcpu=gfx904 %s 2>&1 | FileCheck --check-prefix=GFX904 %s
// RUN: %clang -### -target amdgcn -mcpu=gfx906 %s 2>&1 | FileCheck --check-prefix=GFX906 %s
// GFX600: "-target-cpu" "gfx600"
// TAHITI: "-target-cpu" "tahiti"
@ -113,3 +115,5 @@
// STONEY: "-target-cpu" "stoney"
// GFX900: "-target-cpu" "gfx900"
// GFX902: "-target-cpu" "gfx902"
// GFX904: "-target-cpu" "gfx904"
// GFX906: "-target-cpu" "gfx906"