Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -14,6 +14,26 @@
 
 // The format of this database matches clang/Basic/Builtins.def.
 
+// Taken from ARM. TargetBuiltins.h includes these files without defining
+// LANGBUILTIN
+#if defined(BUILTIN) && !defined(LANGBUILTIN)
+#   define LANGBUILTIN(ID, TYPE, ATTRS, BUILTIN_LANG) BUILTIN(ID, TYPE, ATTRS)
+#endif
+
+//===----------------------------------------------------------------------===//
+// SI+ only builtins.
+//===----------------------------------------------------------------------===//
+
+BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*1", "nc")
+
+LANGBUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_amdgcn_workgroup_id_z, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc", OCLC_LANG)
+
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
@@ -53,6 +73,36 @@
 BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n")
 
 //===----------------------------------------------------------------------===//
+// R600-NI only builtins.
+//===----------------------------------------------------------------------===//
+
+LANGBUILTIN(__builtin_r600_read_workdim, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_r600_read_ngroups_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_ngroups_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_ngroups_z, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_r600_read_global_size_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_global_size_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_global_size_z, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_r600_read_local_size_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_local_size_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_local_size_z, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_r600_read_tgid_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_tgid_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_tgid_z, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_r600_read_tidig_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_tidig_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_tidig_z, "Ui", "nc", OCLC_LANG)
+
+LANGBUILTIN(__builtin_r600_read_global_offset_x, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_global_offset_y, "Ui", "nc", OCLC_LANG)
+LANGBUILTIN(__builtin_r600_read_global_offset_z, "Ui", "nc", OCLC_LANG)
+
+//===----------------------------------------------------------------------===//
 // Legacy names with amdgpu prefix
 //===----------------------------------------------------------------------===//
 
@@ -62,3 +112,4 @@
 BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
 
 #undef BUILTIN
+#undef LANGBUILTIN
Index: lib/Basic/Targets.cpp
===================================================================
--- lib/Basic/Targets.cpp
+++ lib/Basic/Targets.cpp
@@ -2146,6 +2146,8 @@
 const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
+#define LANGBUILTIN(ID, TYPE, ATTRS, LANG)                \
+  { #ID, TYPE, ATTRS, nullptr, LANG, nullptr },
 #include "clang/Basic/BuiltinsAMDGPU.def"
 };
 const char * const AMDGPUTargetInfo::GCCRegNames[] = {
Index: test/CodeGenOpenCL/builtins-amdgcn.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn.cl
+++ test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -275,3 +275,31 @@
 {
   *out = __builtin_amdgpu_ldexp(a, b);
 }
+
+// CHECK-LABEL: @test_get_group_id(
+// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
+// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
+// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()
+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: @test_get_local_id(
+// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x()
+// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y()
+// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z()
+void test_get_local_id(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
+	case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
+	case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
+	default: *out = 0;
+	}
+}
Index: test/CodeGenOpenCL/builtins-r600.cl
===================================================================
--- test/CodeGenOpenCL/builtins-r600.cl
+++ test/CodeGenOpenCL/builtins-r600.cl
@@ -30,3 +30,94 @@
 {
   *out = __builtin_amdgpu_ldexp(a, b);
 }
+
+// CHECK-LABEL: @test_get_work_dim(
+// CHECK: tail call i32 @llvm.r600.read.workdim()
+void test_get_work_dim(global int* out)
+{
+  *out = __builtin_r600_read_workdim();
+}
+
+// CHECK-LABEL: @test_get_num_groups(
+// CHECK: tail call i32 @llvm.r600.read.ngroups.x()
+// CHECK: tail call i32 @llvm.r600.read.ngroups.y()
+// CHECK: tail call i32 @llvm.r600.read.ngroups.z()
+void test_get_num_groups(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_r600_read_ngroups_x(); break;
+	case 1: *out = __builtin_r600_read_ngroups_y(); break;
+	case 2: *out = __builtin_r600_read_ngroups_z(); break;
+	default: *out = 1;
+	}
+}
+
+// CHECK-LABEL: @test_get_global_size(
+// CHECK: tail call i32 @llvm.r600.read.global.size.x()
+// CHECK: tail call i32 @llvm.r600.read.global.size.y()
+// CHECK: tail call i32 @llvm.r600.read.global.size.z()
+void test_get_global_size(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_r600_read_global_size_x(); break;
+	case 1: *out = __builtin_r600_read_global_size_y(); break;
+	case 2: *out = __builtin_r600_read_global_size_z(); break;
+	default: *out = 1;
+	}
+}
+
+// CHECK-LABEL: @test_get_local_size(
+// CHECK: tail call i32 @llvm.r600.read.local.size.x()
+// CHECK: tail call i32 @llvm.r600.read.local.size.y()
+// CHECK: tail call i32 @llvm.r600.read.local.size.z()
+void test_get_local_size(int d, global int * out)
+{
+	switch (d) {
+	case 0: *out = __builtin_r600_read_local_size_x(); break;
+	case 1: *out = __builtin_r600_read_local_size_y(); break;
+	case 2: *out = __builtin_r600_read_local_size_z(); break;
+	default: *out = 1;
+	}
+}
+
+// CHECK-LABEL: @test_get_group_id(
+// CHECK: tail call i32 @llvm.r600.read.tgid.x()
+// CHECK: tail call i32 @llvm.r600.read.tgid.y()
+// CHECK: tail call i32 @llvm.r600.read.tgid.z()
+void test_get_group_id(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_r600_read_tgid_x(); break;
+	case 1: *out = __builtin_r600_read_tgid_y(); break;
+	case 2: *out = __builtin_r600_read_tgid_z(); break;
+	default: *out = 0;
+	}
+}
+
+// CHECK-LABEL: @test_get_local_id(
+// CHECK: tail call i32 @llvm.r600.read.tidig.x()
+// CHECK: tail call i32 @llvm.r600.read.tidig.y()
+// CHECK: tail call i32 @llvm.r600.read.tidig.z()
+void test_get_local_id(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_r600_read_tidig_x(); break;
+	case 1: *out = __builtin_r600_read_tidig_y(); break;
+	case 2: *out = __builtin_r600_read_tidig_z(); break;
+	default: *out = 0;
+	}
+}
+
+// CHECK-LABEL: @test_get_global_offset(
+// CHECK: tail call i32 @llvm.r600.read.global.offset.x()
+// CHECK: tail call i32 @llvm.r600.read.global.offset.y()
+// CHECK: tail call i32 @llvm.r600.read.global.offset.z()
+int test_get_global_offset(int d, global int *out)
+{
+	switch (d) {
+	case 0: *out = __builtin_r600_read_global_offset_x(); break;
+	case 1: *out = __builtin_r600_read_global_offset_y(); break;
+	case 2: *out = __builtin_r600_read_global_offset_z(); break;
+	default: *out = 0;
+	}
+}