diff options
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 28 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error.cl | 20 |
3 files changed, 50 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index d57be9a6d12..15482775488 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -37,6 +37,8 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc") //===----------------------------------------------------------------------===// 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") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 7b4ef08ce30..e33c3ca5266 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -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() diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl index 5e45f7e76c2..6cb2f388225 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -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}} |
