commit 5495f781656011f80e728fcfe01e25aac1701aa2
author Matt Arsenault <Matthew.Arsenault@amd.com> Mon Jun 24 23:34:06 2019 +0000
committer Matt Arsenault <Matthew.Arsenault@amd.com> Mon Jun 24 23:34:06 2019 +0000
tree 29b5e6f072a4b2ed6f0252245ba4ef7843548308
parent e75b1efa39673143af4a2a040179cdb55e1821e7

AMDGPU: Fix missing declaration for mbcnt builtins

llvm-svn: 364251

clang/include/clang/Basic/BuiltinsAMDGPU.def

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 49002f0..e882d3b 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -33,6 +33,9 @@
 BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
 
+BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi", "nc")
+
 //===----------------------------------------------------------------------===//
 // Instruction builtins.
 //===----------------------------------------------------------------------===//

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 3b710a8..e4c40d9 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -578,6 +578,18 @@
   __builtin_amdgcn_ds_gws_sema_p(id);
 }
 
+// CHECK-LABEL: @test_mbcnt_lo(
+// CHECK: call i32 @llvm.amdgcn.mbcnt.lo(i32 %src0, i32 %src1)
+kernel void test_mbcnt_lo(global uint* out, uint src0, uint src1) {
+  *out = __builtin_amdgcn_mbcnt_lo(src0, src1);
+}
+
+// CHECK-LABEL: @test_mbcnt_hi(
+// CHECK: call i32 @llvm.amdgcn.mbcnt.hi(i32 %src0, i32 %src1)
+kernel void test_mbcnt_hi(global uint* out, uint src0, uint src1) {
+  *out = __builtin_amdgcn_mbcnt_hi(src0, src1);
+}
+
 // CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
 // CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }