Author: jvesely Date: Fri Feb 24 22:20:24 2017 New Revision: 296241 URL: http://llvm.org/viewvc/llvm-project?rev=296241&view=rev Log: AMDGPU: export s_sendmsg{halt} instrinsics
Differential Revision: https://reviews.llvm.org/D30366 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=296241&r1=296240&r2=296241&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Feb 24 22:20:24 2017 @@ -37,6 +37,8 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, //===----------------------------------------------------------------------===// BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n") BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n") +BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n") +BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n") Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=296241&r1=296240&r2=296241&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Feb 24 22:20:24 2017 @@ -284,6 +284,34 @@ void test_s_waitcnt() __builtin_amdgcn_s_waitcnt(0); } +// CHECK-LABEL: @test_s_sendmsg +// CHECK: call void @llvm.amdgcn.s.sendmsg( +void test_s_sendmsg() +{ + __builtin_amdgcn_s_sendmsg(1, 0); +} + +// CHECK-LABEL: @test_s_sendmsg_var +// CHECK: call void @llvm.amdgcn.s.sendmsg( +void test_s_sendmsg_var(int in) +{ + __builtin_amdgcn_s_sendmsg(1, in); +} + +// CHECK-LABEL: @test_s_sendmsghalt +// CHECK: call void @llvm.amdgcn.s.sendmsghalt( +void test_s_sendmsghalt() +{ + __builtin_amdgcn_s_sendmsghalt(1, 0); +} + +// CHECK-LABEL: @test_s_sendmsghalt +// CHECK: call void @llvm.amdgcn.s.sendmsghalt( +void test_s_sendmsghalt_var(int in) +{ + __builtin_amdgcn_s_sendmsghalt(1, in); +} + // CHECK-LABEL: @test_s_barrier // CHECK: call void @llvm.amdgcn.s.barrier( void test_s_barrier() Modified: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl?rev=296241&r1=296240&r2=296241&view=diff ============================================================================== --- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl (original) +++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Fri Feb 24 22:20:24 2017 @@ -23,6 +23,26 @@ void test_s_waitcnt(int x) __builtin_amdgcn_s_waitcnt(x); // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}} } +void test_s_sendmsg(int in) +{ + __builtin_amdgcn_s_sendmsg(in, 1); // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}} +} + +void test_s_sendmsg_var(int in1, int in2) +{ + __builtin_amdgcn_s_sendmsg(in1, in2); // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}} +} + +void test_s_sendmsghalt(int in) +{ + __builtin_amdgcn_s_sendmsghalt(in, 1); // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}} +} + +void test_s_sendmsghalt_var(int in1, int in2) +{ + __builtin_amdgcn_s_sendmsghalt(in1, in2); // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}} +} + void test_s_incperflevel(int x) { __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
