This is an archive of the discontinued LLVM Phabricator instance.

Differential D63649

AMDGPU: Fix target builtins for gfx10
ClosedPublic

Authored by arsenm on Jun 21 2019, 6:48 AM.

Details

Reviewers
rampitec
Summary

This wasn't setting some of the features from older generations.

Diff Detail

Event Timeline

arsenm created this revision.Jun 21 2019, 6:48 AM
Herald added subscribers: t-tye, tpr, dstuttard and 5 others. · View Herald TranscriptJun 21 2019, 6:48 AM
rampitec accepted this revision.Jun 21 2019, 9:20 AM

LGTM

This revision is now accepted and ready to land.Jun 21 2019, 9:20 AM
arsenm closed this revision.Jun 21 2019, 6:30 PM

r364123

Revision Contents

Diff 205992

lib/Basic/Targets/AMDGPU.cpp

Show First 20 Lines • Show All 138 Lines • ▼ Show 20 Lines if (isAMDGCN(getTriple())) {
case GK_GFX1011: case GK_GFX1011:
Features["dot1-insts"] = true; Features["dot1-insts"] = true;
Features["dot2-insts"] = true; Features["dot2-insts"] = true;
Features["dot5-insts"] = true; Features["dot5-insts"] = true;
Features["dot6-insts"] = true; Features["dot6-insts"] = true;
LLVM_FALLTHROUGH; LLVM_FALLTHROUGH;
case GK_GFX1010: case GK_GFX1010:
Features["dl-insts"] = true; Features["dl-insts"] = true;
Features["ci-insts"] = true;
Features["16-bit-insts"] = true; Features["16-bit-insts"] = true;
Features["dpp"] = true; Features["dpp"] = true;
Features["gfx8-insts"] = true;
Features["gfx9-insts"] = true; Features["gfx9-insts"] = true;
Features["gfx10-insts"] = true; Features["gfx10-insts"] = true;
Features["s-memrealtime"] = true; Features["s-memrealtime"] = true;
break; break;
case GK_GFX906: case GK_GFX906:
Features["dl-insts"] = true; Features["dl-insts"] = true;
Features["dot1-insts"] = true; Features["dot1-insts"] = true;
Features["dot2-insts"] = true; Features["dot2-insts"] = true;
▲ Show 20 Lines • Show All 186 Lines • Show Last 20 Lines

test/CodeGenOpenCL/amdgpu-features.cl

Show All 9 Lines
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1012 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX1012 %s
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s
// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s
// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" // GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime" // GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+gfx9-insts,+s-memrealtime"
// GFX1010: "target-features"="+16-bit-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx9-insts,+s-memrealtime" // GFX1010: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
// GFX1011: "target-features"="+16-bit-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx9-insts,+s-memrealtime" // GFX1011: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
// GFX1012: "target-features"="+16-bit-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx9-insts,+s-memrealtime" // GFX1012: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot5-insts,+dot6-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx10-insts,+gfx8-insts,+gfx9-insts,+s-memrealtime"
// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime" // GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx8-insts,+s-memrealtime"
// GFX700: "target-features"="+ci-insts,+fp64-fp16-denormals,-fp32-denormals" // GFX700: "target-features"="+ci-insts,+fp64-fp16-denormals,-fp32-denormals"
// GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals" // GFX600: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
// GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals" // GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals"
kernel void test() {} kernel void test() {}

test/CodeGenOpenCL/builtins-amdgcn-ci.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
typedef unsigned int uint; typedef unsigned int uint;
// CHECK-LABEL: @test_s_dcache_inv_vol // CHECK-LABEL: @test_s_dcache_inv_vol
// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol( // CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
void test_s_dcache_inv_vol() void test_s_dcache_inv_vol()
{ {
__builtin_amdgcn_s_dcache_inv_vol(); __builtin_amdgcn_s_dcache_inv_vol();
Show All 15 Lines

test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err-clamp.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -verify -S -emit-llvm -o - %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -verify -S -emit-llvm -o - %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -verify -S -emit-llvm -o - %s
typedef unsigned int uint; typedef unsigned int uint;
typedef half __attribute__((ext_vector_type(2))) half2; typedef half __attribute__((ext_vector_type(2))) half2;
typedef short __attribute__((ext_vector_type(2))) short2; typedef short __attribute__((ext_vector_type(2))) short2;
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2; typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
kernel void builtins_amdgcn_dl_insts_err( kernel void builtins_amdgcn_dl_insts_err(
global float *fOut, global int *siOut, global uint *uiOut, global float *fOut, global int *siOut, global uint *uiOut,
Show All 14 Lines

test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -S -emit-llvm -o - %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -verify -S -emit-llvm -o - %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -verify -S -emit-llvm -o - %s
typedef unsigned int uint; typedef unsigned int uint;
typedef half __attribute__((ext_vector_type(2))) half2; typedef half __attribute__((ext_vector_type(2))) half2;
typedef short __attribute__((ext_vector_type(2))) short2; typedef short __attribute__((ext_vector_type(2))) short2;
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2; typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
kernel void builtins_amdgcn_dl_insts_err( kernel void builtins_amdgcn_dl_insts_err(
global float *fOut, global int *siOut, global uint *uiOut, global float *fOut, global int *siOut, global uint *uiOut,
Show All 24 Lines

test/CodeGenOpenCL/builtins-amdgcn-dl-insts.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
typedef unsigned int uint; typedef unsigned int uint;
typedef half __attribute__((ext_vector_type(2))) half2; typedef half __attribute__((ext_vector_type(2))) half2;
typedef short __attribute__((ext_vector_type(2))) short2; typedef short __attribute__((ext_vector_type(2))) short2;
typedef unsigned short __attribute__((ext_vector_type(2))) ushort2; typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
// CHECK-LABEL: @builtins_amdgcn_dl_insts // CHECK-LABEL: @builtins_amdgcn_dl_insts
// CHECK: call float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 false) // CHECK: call float @llvm.amdgcn.fdot2(<2 x half> %v2hA, <2 x half> %v2hB, float %fC, i1 false)
▲ Show 20 Lines • Show All 45 Lines • Show Last 20 Lines

test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
#pragma OPENCL EXTENSION cl_khr_fp16 : enable #pragma OPENCL EXTENSION cl_khr_fp16 : enable
// CHECK-LABEL: @test_fmed3_f16 // CHECK-LABEL: @test_fmed3_f16
// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c) // CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c)
void test_fmed3_f16(global half* out, half a, half b, half c) void test_fmed3_f16(global half* out, half a, half b, half c)
{ {
*out = __builtin_amdgcn_fmed3h(a, b, c); *out = __builtin_amdgcn_fmed3h(a, b, c);
} }

test/CodeGenOpenCL/builtins-amdgcn-vi.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
#pragma OPENCL EXTENSION cl_khr_fp16 : enable #pragma OPENCL EXTENSION cl_khr_fp16 : enable
typedef unsigned long ulong; typedef unsigned long ulong;
// CHECK-LABEL: @test_div_fixup_f16 // CHECK-LABEL: @test_div_fixup_f16
// CHECK: call half @llvm.amdgcn.div.fixup.f16 // CHECK: call half @llvm.amdgcn.div.fixup.f16
void test_div_fixup_f16(global half* out, half a, half b, half c) void test_div_fixup_f16(global half* out, half a, half b, half c)
▲ Show 20 Lines • Show All 112 Lines • Show Last 20 Lines