From 780813dafe4b56c8f845cf3f356b066d62d1ac49 Mon Sep 17 00:00:00 2001
From: Jan Vesely <jan.vesely@rutgers.edu>
Date: Sat, 25 Feb 2017 04:20:24 +0000
Subject: [PATCH] AMDGPU: export s_sendmsg{halt} instrinsics

Differential Revision: https://reviews.llvm.org/D30366

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@296241 91177308-0d34-0410-b5e6-96231b3b80d8
---
 include/clang/Basic/BuiltinsAMDGPU.def   |  2 ++
 test/CodeGenOpenCL/builtins-amdgcn.cl    | 28 ++++++++++++++++++++++++
 test/SemaOpenCL/builtins-amdgcn-error.cl | 20 +++++++++++++++++
 3 files changed, 50 insertions(+)

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index d57be9a6d1..1548277548 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/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/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index 7b4ef08ce3..e33c3ca526 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/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/test/SemaOpenCL/builtins-amdgcn-error.cl b/test/SemaOpenCL/builtins-amdgcn-error.cl
index 5e45f7e76c..6cb2f38822 100644
--- a/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/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}}
-- 
2.40.0