Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -132,6 +132,8 @@
 TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
 TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
 TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_release_all, "vUi", "n", "ci-insts")
+TARGET_BUILTIN(__builtin_amdgcn_is_local, "bvC*0", "nc", "flat-address-space")
+TARGET_BUILTIN(__builtin_amdgcn_is_private, "bvC*0", "nc", "flat-address-space")
 
 //===----------------------------------------------------------------------===//
 // Interpolation builtins.
Index: lib/Basic/Targets/AMDGPU.cpp
===================================================================
--- lib/Basic/Targets/AMDGPU.cpp
+++ lib/Basic/Targets/AMDGPU.cpp
@@ -145,6 +145,7 @@
     case GK_GFX1010:
       Features["dl-insts"] = true;
       Features["ci-insts"] = true;
+      Features["flat-address-space"] = true;
       Features["16-bit-insts"] = true;
       Features["dpp"] = true;
       Features["gfx8-insts"] = true;
@@ -184,6 +185,7 @@
     case GK_GFX701:
     case GK_GFX700:
       Features["ci-insts"] = true;
+      Features["flat-address-space"] = true;
       LLVM_FALLTHROUGH;
     case GK_GFX601:
     case GK_GFX600:
Index: test/CodeGenOpenCL/builtins-amdgcn-ci.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -1,8 +1,8 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
 
 typedef unsigned int uint;
 
@@ -21,8 +21,36 @@
 }
 
 // CHECK-LABEL: @test_gws_sema_release_all(
-// CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %id)
+// CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %{{[0-9]+}})
 void test_gws_sema_release_all(uint id)
 {
   __builtin_amdgcn_ds_gws_sema_release_all(id);
 }
+
+// CHECK-LABEL: @test_is_local(
+// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8*
+// CHECK: call i1 @llvm.amdgcn.is.local(i8* [[CAST]]
+int test_is_local(const int* ptr) {
+  return __builtin_amdgcn_is_local(ptr);
+}
+
+// CHECK-LABEL: @test_is_private(
+// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8*
+// CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]]
+int test_is_private(const int* ptr) {
+  return __builtin_amdgcn_is_private(ptr);
+}
+
+// CHECK-LABEL: @test_is_local_global(
+// CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8*
+// CHECK: call i1 @llvm.amdgcn.is.local(i8* [[CAST]]
+int test_is_local_global(const global int* ptr) {
+  return __builtin_amdgcn_is_local(ptr);
+}
+
+// CHECK-LABEL: @test_is_private_global(
+// CHECK: [[CAST:%[0-9]+]] = addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i8*
+// CHECK: call i1 @llvm.amdgcn.is.private(i8* [[CAST]]
+int test_is_private_global(const global int* ptr) {
+  return __builtin_amdgcn_is_private(ptr);
+}
Index: test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
===================================================================
--- /dev/null
+++ test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+
+void test_flat_address_space_builtins(int* ptr)
+{
+  (void)__builtin_amdgcn_is_local(ptr); // expected-error {{'__builtin_amdgcn_is_local' needs target feature flat-address-space}}
+  (void)__builtin_amdgcn_is_private(ptr); // expected-error {{'__builtin_amdgcn_is_private' needs target feature flat-address-space}}
+}