Author: arsenm
Date: Mon Jun 24 16:34:06 2019
New Revision: 364251
URL: http://llvm.org/viewvc/llvm-project?rev=364251&view=rev
Log:
AMDGPU: Fix missing declaration for mbcnt builtins
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=364251&r1=364250&r2=364251&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Mon Jun 24 16:34:06 2019
@@ -33,6 +33,9 @@ BUILTIN(__builtin_amdgcn_workitem_id_x,
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.
//===----------------------------------------------------------------------===//
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=364251&r1=364250&r2=364251&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Mon Jun 24 16:34:06 2019
@@ -578,6 +578,18 @@ kernel void test_gws_sema_p(uint id) {
__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 }
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
