AMDGPU: Add __builtin_amdgcn_workgroup_count_* builtins

include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workgroup_id_z, "Ui", "nc")
 
 BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc")
 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.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
 BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")

lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
 
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec"));
     CI->setConvergent();
     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);
   case AMDGPU::BI__builtin_amdgcn_workitem_id_y:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_y, 0, 1024);
   case AMDGPU::BI__builtin_amdgcn_workitem_id_z:
     return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 1024);
 

test/CodeGenOpenCL/builtins-amdgcn.cl

void test_get_group_id(int d, global int *out)
   switch (d) {
   case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;
   case 1: *out = __builtin_amdgcn_workgroup_id_y(); break;
   case 2: *out = __builtin_amdgcn_workgroup_id_z(); break;
   default: *out = 0;
   }
 }
 
+// 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]]
 // CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[WI_RANGE]]
 void test_get_local_id(int d, global int *out)
 {
   switch (d) {
   case 0: *out = __builtin_amdgcn_workitem_id_x(); break;