[llvm] e630d9b - AMDGPU/clang: Remove target features from address space test builtins
Matt Arsenault via llvm-commits
llvm-commits at lists.llvm.org
Thu Dec 29 15:46:46 PST 2022
Author: Matt Arsenault
Date: 2022-12-29T18:46:41-05:00
New Revision: e630d9b299822810bba8f3d0457004d1b4c39bef
URL: https://github.com/llvm/llvm-project/commit/e630d9b299822810bba8f3d0457004d1b4c39bef
DIFF: https://github.com/llvm/llvm-project/commit/e630d9b299822810bba8f3d0457004d1b4c39bef.diff
LOG: AMDGPU/clang: Remove target features from address space test builtins
It turns out we can codegen these on targets without flat addressing,
although the runtime probably didn't put anything useful there. The
proper diagnostic would be to disallow flat pointer uses or languages
with them, not this one edge case. Allows removing one of the special
cases requiring subtarget support in the device libraries.
Added:
clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl
Modified:
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.private.ll
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.shared.ll
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 9e717099f7777..5c39bae0d5d46 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -164,14 +164,18 @@ BUILTIN(__builtin_amdgcn_sicmpl, "WUiWiWiIi", "nc")
BUILTIN(__builtin_amdgcn_fcmp, "WUiddIi", "nc")
BUILTIN(__builtin_amdgcn_fcmpf, "WUiffIi", "nc")
+//===----------------------------------------------------------------------===//
+// Flat addressing builtins.
+//===----------------------------------------------------------------------===//
+BUILTIN(__builtin_amdgcn_is_shared, "bvC*0", "nc")
+BUILTIN(__builtin_amdgcn_is_private, "bvC*0", "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.
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl
new file mode 100644
index 0000000000000..57ee1c1a64711
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-flat-address-space.cl
@@ -0,0 +1,22 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tahiti -S -emit-llvm -disable-llvm-passes -o - %s | FileCheck -enable-var-scope %s
+
+// SI did not actually support flat addressing, but we can codegen the address
+// space test builtins. The target specfic part is a load from the implicit
+// argument buffer to use for the high pointer bits. It's just that buffer won't
+// be initialized to something useful. The proper way to diagnose invalid flat
+// usage is to forbid flat pointers on unsupported targets.
+
+// CHECK-LABEL: @test_is_shared_global(
+// CHECK: [[CAST:%[0-9]+]] = addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr
+// CHECK: call i1 @llvm.amdgcn.is.shared(ptr [[CAST]]
+int test_is_shared_global(const global int* ptr) {
+ return __builtin_amdgcn_is_shared(ptr);
+}
+
+// CHECK-LABEL: @test_is_private_global(
+// CHECK: [[CAST:%[0-9]+]] = addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr
+// CHECK: call i1 @llvm.amdgcn.is.private(ptr [[CAST]]
+int test_is_private_global(const global int* ptr) {
+ return __builtin_amdgcn_is_private(ptr);
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index eb1fa22e7fac2..9a20dba7490b4 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -1,5 +1,5 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
+// RUN: %clang_cc1 -no-opaque-pointers -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tahiti -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
index 0d759bac74a72..86ee55a646de8 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-flat-address-space.cl
@@ -1,8 +1,12 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+// expected-no-diagnostics
+
+// Make sure no warning is produced on due to dead "flat-address-space" feature.
+__attribute__((target("flat-address-space")))
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}}
+ (void)__builtin_amdgcn_is_shared(ptr);
+ (void)__builtin_amdgcn_is_private(ptr);
}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.private.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.private.ll
index 2a92101203ebe..ff58d05f699e7 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.private.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.private.ll
@@ -1,8 +1,9 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,CI %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=hawaii -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,CI %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX9 %s
; GCN-LABEL: {{^}}is_private_vgpr:
-; GCN-DAG: {{flat|global}}_load_dwordx2 v{{\[[0-9]+}}:[[PTR_HI:[0-9]+]]]
+; GCN-DAG: {{flat|global|buffer}}_load_dwordx2 v{{\[[0-9]+}}:[[PTR_HI:[0-9]+]]]
; CI-DAG: s_load_dword [[APERTURE:s[0-9]+]], s[4:5], 0x11
; CI: v_cmp_eq_u32_e32 vcc, [[APERTURE]], v[[PTR_HI]]
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.shared.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.shared.ll
index 1a10a9c15dc53..7479fc87e3188 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.shared.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.is.shared.ll
@@ -1,8 +1,9 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,CI %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=hawaii -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,CI %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX9 %s
; GCN-LABEL: {{^}}is_local_vgpr:
-; GCN-DAG: {{flat|global}}_load_dwordx2 v{{\[[0-9]+}}:[[PTR_HI:[0-9]+]]]
+; GCN-DAG: {{flat|global|buffer}}_load_dwordx2 v{{\[[0-9]+}}:[[PTR_HI:[0-9]+]]]
; CI-DAG: s_load_dword [[APERTURE:s[0-9]+]], s[4:5], 0x10
; GFX9: s_mov_b64 s[{{[0-9]+}}:[[HI:[0-9]+]]], src_shared_base
More information about the llvm-commits
mailing list