This is an archive of the discontinued LLVM Phabricator instance.

Differential D60293

[AMDGPU] rename vi-insts into gfx8-insts
ClosedPublic

Authored by rampitec on Apr 4 2019, 4:31 PM.

Details

Reviewers
kzhuravl
msearles
Commits
rG1d9f286ecb81: [AMDGPU] rename vi-insts into gfx8-insts
rC357792: [AMDGPU] rename vi-insts into gfx8-insts
rL357792: [AMDGPU] rename vi-insts into gfx8-insts

Diff Detail

Repository
rL LLVM

Event Timeline

rampitec created this revision.Apr 4 2019, 4:31 PM
Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptApr 4 2019, 4:31 PM
msearles accepted this revision.Apr 4 2019, 4:39 PM

LGTM

This revision is now accepted and ready to land.Apr 4 2019, 4:39 PM
Closed by commit rL357792: [AMDGPU] rename vi-insts into gfx8-insts (authored by rampitec). · Explain WhyApr 5 2019, 11:23 AM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptApr 5 2019, 11:23 AM

Revision Contents

PathSize
cfe/
trunk/
include/
clang/
Basic/
2 lines
lib/
Basic/
Targets/
2 lines
test/
CodeGenOpenCL/
6 lines
SemaOpenCL/
2 lines

Diff 193916

cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def

Show First 20 Lines • Show All 127 Lines • ▼ Show 20 Lines
TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime") TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp") TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp")
TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nc", "dpp") TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb", "nc", "dpp")
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "vi-insts") TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "gfx8-insts")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// GFX9+ only builtins. // GFX9+ only builtins.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts") TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
Show All 37 Lines

cfe/trunk/lib/Basic/Targets/AMDGPU.cpp

Show First 20 Lines • Show All 144 Lines • ▼ Show 20 Lines if (isAMDGCN(getTriple())) {
case GK_GFX902: case GK_GFX902:
case GK_GFX900: case GK_GFX900:
Features["gfx9-insts"] = true; Features["gfx9-insts"] = true;
LLVM_FALLTHROUGH; LLVM_FALLTHROUGH;
case GK_GFX810: case GK_GFX810:
case GK_GFX803: case GK_GFX803:
case GK_GFX802: case GK_GFX802:
case GK_GFX801: case GK_GFX801:
Features["vi-insts"] = true; Features["gfx8-insts"] = true;
Features["16-bit-insts"] = true; Features["16-bit-insts"] = true;
Features["dpp"] = true; Features["dpp"] = true;
Features["s-memrealtime"] = true; Features["s-memrealtime"] = true;
LLVM_FALLTHROUGH; LLVM_FALLTHROUGH;
case GK_GFX704: case GK_GFX704:
case GK_GFX703: case GK_GFX703:
case GK_GFX702: case GK_GFX702:
case GK_GFX701: case GK_GFX701:
▲ Show 20 Lines • Show All 163 Lines • Show Last 20 Lines

cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// Check that appropriate features are defined for every supported AMDGPU // Check that appropriate features are defined for every supported AMDGPU
// "-target" and "-mcpu" options. // "-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 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 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %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,+gfx9-insts,+s-memrealtime,+vi-insts" // 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,+gfx9-insts,+s-memrealtime,+vi-insts" // 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"
// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+s-memrealtime,+vi-insts" // 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() {}

cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-vi.cl

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s // RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s
void test_vi_s_dcache_wb() void test_vi_s_dcache_wb()
{ {
__builtin_amdgcn_s_dcache_wb(); // expected-error {{'__builtin_amdgcn_s_dcache_wb' needs target feature vi-insts}} __builtin_amdgcn_s_dcache_wb(); // expected-error {{'__builtin_amdgcn_s_dcache_wb' needs target feature gfx8-insts}}
} }