[PATCH] D22790: AMDGPU: Add __builtin_amdgcn_workgroup_count_* builtins

[Tom Stellard via llvm-commits]

tstellarAMD created this revision.
tstellarAMD added a reviewer: arsenm.
tstellarAMD added a subscriber: llvm-commits.
Herald added a subscriber: kzhuravl.

https://reviews.llvm.org/D22790

Files:
  include/clang/Basic/BuiltinsAMDGPU.def
  lib/CodeGen/CGBuiltin.cpp
  test/CodeGenOpenCL/builtins-amdgcn.cl

Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -296,6 +296,19 @@
 	}
 }
 
+// CHECK-LABEL: @get_num_groups(
+// CHECK: call i32 @llvm.amdgcn.workgroup.count.x()
+// CHECK: call i32 @llvm.amdgcn.workgroup.count.y()
+// CHECK: call i32 @llvm.amdgcn.workgroup.count.z()
+void get_num_groups(int dim, global int *out) {
+  switch (dim) {
+  case 0: *out = __builtin_amdgcn_workgroup_count_x(); break;
+  case 1: *out = __builtin_amdgcn_workgroup_count_y(); break;
+  case 2: *out = __builtin_amdgcn_workgroup_count_z(); break;
+  default: *out = 0; break;
+  }
+}
+
 // CHECK-LABEL: @test_get_local_id(
 // CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]]
 // CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]]
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -7672,6 +7672,14 @@
     return CI;
   }
 
+  // amdgcn workgroup count
+  case AMDGPU::BI__builtin_amdgcn_workgroup_count_x:
+    return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_workgroup_count_x);
+  case AMDGPU::BI__builtin_amdgcn_workgroup_count_y:
+    return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_workgroup_count_y);
+  case AMDGPU::BI__builtin_amdgcn_workgroup_count_z:
+    return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_workgroup_count_z);
+
   // amdgcn workitem
   case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0, 1024);
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -32,6 +32,10 @@
 BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
 
+BUILTIN(__builtin_amdgcn_workgroup_count_x, "Ui", "nc")
+BUILTIN(__builtin_amdgcn_workgroup_count_y, "Ui", "nc")
+BUILTIN(__builtin_amdgcn_workgroup_count_z, "Ui", "nc")
+
 //===----------------------------------------------------------------------===//
 // Instruction builtins.
 //===----------------------------------------------------------------------===//


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D22790.65448.patch
Type: text/x-patch
Size: 2386 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160726/fd714b84/attachment.bin>