AMDGPU: Add builtins for is_local/is_private

[b-sumner]

Looks fine to me. Thanks!

[arsenm]

rename

[b-sumner]

Looks fine to me.

[rampitec]

LGTM

[rampitec]

Didn't you forget to update test/CodeGenOpenCL/amdgpu-features.cl?

[arsenm]

r371010

include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiLUiUiV4Ui", "nc")
 
 //===----------------------------------------------------------------------===//
 // CI+ only builtins.
 //===----------------------------------------------------------------------===//
 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_shared, "bvC*0", "nc", "flat-address-space")
+TARGET_BUILTIN(__builtin_amdgcn_is_private, "bvC*0", "nc", "flat-address-space")
 
 //===----------------------------------------------------------------------===//
 // Interpolation builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_interp_p1_f16, "ffUiUibUi", "nc")
 BUILTIN(__builtin_amdgcn_interp_p2_f16, "hffUiUibUi", "nc")
 BUILTIN(__builtin_amdgcn_interp_p1, "ffUiUiUi", "nc")
 BUILTIN(__builtin_amdgcn_interp_p2, "fffUiUiUi", "nc")

lib/Basic/Targets/AMDGPU.cpp

case GK_GFX1011:
       Features["dot1-insts"] = true;
       Features["dot2-insts"] = true;
       Features["dot5-insts"] = true;
       Features["dot6-insts"] = true;
       LLVM_FALLTHROUGH;
     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;
       Features["gfx9-insts"] = true;
       Features["gfx10-insts"] = true;
       Features["s-memrealtime"] = true;
       break;
     case GK_GFX908:
case GK_GFX801:
       Features["s-memrealtime"] = true;
       LLVM_FALLTHROUGH;
     case GK_GFX704:
     case GK_GFX703:
     case GK_GFX702:
     case GK_GFX701:
     case GK_GFX700:
       Features["ci-insts"] = true;
+      Features["flat-address-space"] = true;
       LLVM_FALLTHROUGH;
     case GK_GFX601:
     case GK_GFX600:
       break;
     case GK_NONE:
       return false;
     default:
       llvm_unreachable("Unhandled GPU!");

test/CodeGenOpenCL/builtins-amdgcn-ci.cl

 // 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;
 
 // CHECK-LABEL: @test_s_dcache_inv_vol
 // CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
 void test_s_dcache_inv_vol()
 {
   __builtin_amdgcn_s_dcache_inv_vol();
 }
 
 // CHECK-LABEL: @test_buffer_wbinvl1_vol
 // CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol()
 void test_buffer_wbinvl1_vol()
 {
   __builtin_amdgcn_buffer_wbinvl1_vol();
 }
 
 // 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_shared(
+// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8*
+// CHECK: call i1 @llvm.amdgcn.is.local(i8* [[CAST]]
+int test_is_shared(const int* ptr) {
+  return __builtin_amdgcn_is_shared(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);
+}

test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl

+// 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_shared(ptr); // expected-error {{'__builtin_amdgcn_is_shared' needs target feature flat-address-space}}
+  (void)__builtin_amdgcn_is_private(ptr); // expected-error {{'__builtin_amdgcn_is_private' needs target feature flat-address-space}}
+}