Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-01-25 10:48:14 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

AMDGPU: Add another missing builtin

Browse Source 
llvm-svn: 339395
This commit is contained in:
Matt Arsenault 2018-08-09 22:18:37 +00:00
parent 941095100f
commit 94abc57e37
3 changed files with 10 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -104,6 +104,7 @@ BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
// CI+ only builtins.
//===----------------------------------------------------------------------===//
TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
//===----------------------------------------------------------------------===//
// VI+ only builtins.

7
clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
View File

@ -10,3 +10,10 @@ void test_s_dcache_inv_vol()
__builtin_amdgcn_s_dcache_inv_vol();
}
// CHECK-LABEL: @test_buffer_wbinvl1_vol
// CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol()
void test_buffer_wbinvl1_vol()
{
__builtin_amdgcn_buffer_wbinvl1_vol();
}

3
clang/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
View File

@ -1,7 +1,8 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
void test_ci_s_dcache_inv_vol()
void test_ci_biltins()
{
__builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
__builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}}
}