AMDGPU: Add Vega12 and Vega20

Changes by
  Matt Arsenault
  Konstantin Zhuravlyov

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@331216 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 46687f75984cf5f230ad19b984bd76090c9c4f0a..21b15314e8ba6a91347e8c543e492af88fc70fc2 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -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.
 //===----------------------------------------------------------------------===//

diff --git a/lib/Basic/Targets/AMDGPU.cpp b/lib/Basic/Targets/AMDGPU.cpp
index d956d132a5b641409e104639ff140d7435f52411..b6b9aa2f1244cbb1f2cd69f50dbc60ec41c3b3a2 100644 (file)
--- a/lib/Basic/Targets/AMDGPU.cpp
+++ b/lib/Basic/Targets/AMDGPU.cpp
@@ -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;

diff --git a/lib/Basic/Targets/AMDGPU.h b/lib/Basic/Targets/AMDGPU.h
index f7c477659dc1a6491c98883f1e8bb166d9c8119f..6a264f8456abe5f67a59b60bb1b4ee1d7bbeb089 100644 (file)
--- a/lib/Basic/Targets/AMDGPU.h
+++ b/lib/Basic/Targets/AMDGPU.h
@@ -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);

diff --git a/test/CodeGenOpenCL/amdgpu-features.cl b/test/CodeGenOpenCL/amdgpu-features.cl
new file mode 100644 (file)
index 0000000..a1815ce
--- /dev/null
+++ b/test/CodeGenOpenCL/amdgpu-features.cl
@@ -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() {}

diff --git a/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl b/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
new file mode 100644 (file)
index 0000000..bfc811e
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
@@ -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}}
+}

diff --git a/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl b/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl
new file mode 100644 (file)
index 0000000..c407bd6
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl
@@ -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);
+}

diff --git a/test/Driver/amdgpu-macros.cl b/test/Driver/amdgpu-macros.cl
index a047a7cee50f0915b87833d34a2600c5ae2188bd..c57d3805be7a7a0b5dae07bad2f41c11cae1242c 100644 (file)
--- a/test/Driver/amdgpu-macros.cl
+++ b/test/Driver/amdgpu-macros.cl
@@ -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

diff --git a/test/Driver/amdgpu-mcpu.cl b/test/Driver/amdgpu-mcpu.cl
index 265d0f7797f5b87365f158840ed1f8556ffc84a4..c036e2ef093a179124256c759d023c100ff4634b 100644 (file)
--- a/test/Driver/amdgpu-mcpu.cl
+++ b/test/Driver/amdgpu-mcpu.cl
@@ -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"