Author: jvesely Date: Fri Feb 24 22:20:22 2017 New Revision: 296240 URL: http://llvm.org/viewvc/llvm-project?rev=296240&view=rev Log: AMDGPU: export l1 cache invalidation intrinsics
Differential Revision: https://reviews.llvm.org/D30360 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=296240&r1=296239&r2=296240&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Feb 24 22:20:22 2017 @@ -39,6 +39,8 @@ BUILTIN(__builtin_amdgcn_s_getreg, "UiIi 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") Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=296240&r1=296239&r2=296240&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Feb 24 22:20:22 2017 @@ -263,6 +263,20 @@ void test_class_f64(global double* out, *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() _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
