r331216 - AMDGPU: Add Vega12 and Vega20
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 30 12:08:27 PDT 2018
Author: arsenm
Date: Mon Apr 30 12:08:27 2018
New Revision: 331216
URL: http://llvm.org/viewvc/llvm-project?rev=331216&view=rev
Log:
AMDGPU: Add Vega12 and Vega20
Changes by
Matt Arsenault
Konstantin Zhuravlyov
Added:
cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
cfe/trunk/lib/Basic/Targets/AMDGPU.h
cfe/trunk/test/Driver/amdgpu-macros.cl
cfe/trunk/test/Driver/amdgpu-mcpu.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=331216&r1=331215&r2=331216&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon Apr 30 12:08:27 2018
@@ -121,6 +121,18 @@ TARGET_BUILTIN(__builtin_amdgcn_mov_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")
Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.cpp?rev=331216&r1=331215&r2=331216&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets/AMDGPU.cpp (original)
+++ cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Mon Apr 30 12:08:27 2018
@@ -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;
Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.h?rev=331216&r1=331215&r2=331216&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets/AMDGPU.h (original)
+++ cfe/trunk/lib/Basic/Targets/AMDGPU.h Mon Apr 30 12:08:27 2018
@@ -78,9 +78,11 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTarg
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 AMDGPUTarg
{{"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 AMDGPUTarg
{{"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);
Added: cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl?rev=331216&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl Mon Apr 30 12:08:27 2018
@@ -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() {}
Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl?rev=331216&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl Mon Apr 30 12:08:27 2018
@@ -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}}
+}
Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl?rev=331216&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl Mon Apr 30 12:08:27 2018
@@ -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);
+}
Modified: cfe/trunk/test/Driver/amdgpu-macros.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Driver/amdgpu-macros.cl?rev=331216&r1=331215&r2=331216&view=diff
==============================================================================
--- cfe/trunk/test/Driver/amdgpu-macros.cl (original)
+++ cfe/trunk/test/Driver/amdgpu-macros.cl Mon Apr 30 12:08:27 2018
@@ -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
Modified: cfe/trunk/test/Driver/amdgpu-mcpu.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Driver/amdgpu-mcpu.cl?rev=331216&r1=331215&r2=331216&view=diff
==============================================================================
--- cfe/trunk/test/Driver/amdgpu-mcpu.cl (original)
+++ cfe/trunk/test/Driver/amdgpu-mcpu.cl Mon Apr 30 12:08:27 2018
@@ -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"
More information about the cfe-commits
mailing list