diff options
| author | Jan Vesely <jan.vesely@rutgers.edu> | 2017-02-25 04:20:22 +0000 |
|---|---|---|
| committer | Jan Vesely <jan.vesely@rutgers.edu> | 2017-02-25 04:20:22 +0000 |
| commit | c255097517131c5c40e3ba408ddf44924ec4d521 (patch) | |
| tree | 74c5619f347885b74287f5e0a4faa63b8b5785be /clang | |
| parent | d26dbb389f1e693a8a7d085eb026c68718f68455 (diff) | |
| download | bcm5719-llvm-c255097517131c5c40e3ba408ddf44924ec4d521.tar.gz bcm5719-llvm-c255097517131c5c40e3ba408ddf44924ec4d521.zip | |
AMDGPU: export l1 cache invalidation intrinsics
Differential Revision: https://reviews.llvm.org/D30360
llvm-svn: 296240
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 14 |
2 files changed, 16 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index f39a24157a6..d57be9a6d12 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -39,6 +39,8 @@ BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n") BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") +BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n") +BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n") BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index e9e765b521b..7b4ef08ce30 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -263,6 +263,20 @@ void test_class_f64(global double* out, double a, int b) *out = __builtin_amdgcn_class(a, b); } +// CHECK-LABEL: @test_buffer_wbinvl1 +// CHECK: call void @llvm.amdgcn.buffer.wbinvl1( +void test_buffer_wbinvl1() +{ + __builtin_amdgcn_buffer_wbinvl1(); +} + +// CHECK-LABEL: @test_s_dcache_inv +// CHECK: call void @llvm.amdgcn.s.dcache.inv( +void test_s_dcache_inv() +{ + __builtin_amdgcn_s_dcache_inv(); +} + // CHECK-LABEL: @test_s_waitcnt // CHECK: call void @llvm.amdgcn.s.waitcnt( void test_s_waitcnt() |
