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")
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;
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 {
{{"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
{{"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);
--- /dev/null
+// 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() {}
--- /dev/null
+// 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}}
+}
--- /dev/null
+// 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);
+}
// 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
// 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
// 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
// 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
// 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
// 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
// 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
// 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"
// STONEY: "-target-cpu" "stoney"
// GFX900: "-target-cpu" "gfx900"
// GFX902: "-target-cpu" "gfx902"
+// GFX904: "-target-cpu" "gfx904"
+// GFX906: "-target-cpu" "gfx906"