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.
 //===----------------------------------------------------------------------===//
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: 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]]