dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

AMDGPU: Add another missing builtin 

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@339395 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Matt Arsenault 2018-08-09 22:18:37 +00:00
parent 9a2dafaebb
commit f2b72fdc01
3 changed files with 10 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

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

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

@ -10,3 +10,10 @@ void test_s_dcache_inv_vol()
__builtin_amdgcn_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
test/SemaOpenCL/builtins-amdgcn-error-ci.cl
View File

@ -1,7 +1,8 @@
// 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
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_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}}
} }