summaryrefslogtreecommitdiffstats
path: root/clang
diff options
context:
space:
mode:
Diffstat (limited to 'clang')
-rw-r--r--clang/include/clang/Basic/BuiltinsAMDGPU.def12
-rw-r--r--clang/lib/Basic/Targets/AMDGPU.cpp4
-rw-r--r--clang/lib/Basic/Targets/AMDGPU.h8
-rw-r--r--clang/test/CodeGenOpenCL/amdgpu-features.cl12
-rw-r--r--clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl25
-rw-r--r--clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl36
-rw-r--r--clang/test/Driver/amdgpu-macros.cl14
-rw-r--r--clang/test/Driver/amdgpu-mcpu.cl4
8 files changed, 113 insertions, 2 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 46687f75984..21b15314e8b 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -121,6 +121,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.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index d956d132a5b..b6b9aa2f124 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/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/clang/lib/Basic/Targets/AMDGPU.h b/clang/lib/Basic/Targets/AMDGPU.h
index f7c477659dc..6a264f8456a 100644
--- a/clang/lib/Basic/Targets/AMDGPU.h
+++ b/clang/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/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl
new file mode 100644
index 00000000000..a1815cedcd4
--- /dev/null
+++ b/clang/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/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
new file mode 100644
index 00000000000..bfc811e2955
--- /dev/null
+++ b/clang/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/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl
new file mode 100644
index 00000000000..c407bd6fe15
--- /dev/null
+++ b/clang/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/clang/test/Driver/amdgpu-macros.cl b/clang/test/Driver/amdgpu-macros.cl
index a047a7cee50..c57d3805be7 100644
--- a/clang/test/Driver/amdgpu-macros.cl
+++ b/clang/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/clang/test/Driver/amdgpu-mcpu.cl b/clang/test/Driver/amdgpu-mcpu.cl
index 265d0f7797f..c036e2ef093 100644
--- a/clang/test/Driver/amdgpu-mcpu.cl
+++ b/clang/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"
OpenPOWER on IntegriCloud