arsenm created this revision.
arsenm added reviewers: rampitec, b-sumner.
Herald added subscribers: t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng,
jvesely, kzhuravl.
arsenm added a parent revision: D66197: AMDGPU: Add intrinsics for address
space identification.
https://reviews.llvm.org/D66198
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/Basic/Targets/AMDGPU.cpp
test/CodeGenOpenCL/builtins-amdgcn-ci.cl
test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
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}}
+}
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: 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: 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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
