From: Matt Arsenault Date: Thu, 9 Aug 2018 22:18:37 +0000 (+0000) Subject: AMDGPU: Add another missing builtin X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=f2b72fdc01b332836a1d31ce3cb70424c64ad178;p=clang AMDGPU: Add another missing builtin git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@339395 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index 73204dc747..35f3a887de 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -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. diff --git a/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/test/CodeGenOpenCL/builtins-amdgcn-ci.cl index 023e761433..41275268db 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -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(); +} + diff --git a/test/SemaOpenCL/builtins-amdgcn-error-ci.cl b/test/SemaOpenCL/builtins-amdgcn-error-ci.cl index 282b45fdb1..2f656582be 100644 --- a/test/SemaOpenCL/builtins-amdgcn-error-ci.cl +++ b/test/SemaOpenCL/builtins-amdgcn-error-ci.cl @@ -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}} }