https://github.com/jmmartinez created
https://github.com/llvm/llvm-project/pull/135376
The vmem-to-lds-loads-insts feature is only available on gfx9/10. While
target-parser was also enabling it for gfx6,7,8.
From 90c95af57317103f399ee4059a4e4ec01ddc98de Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <juama...@amd.com>
Date: Fri, 11 Apr 2025 15:00:44 +0200
Subject: [PATCH] [AMDGPU] Incoherence between TargetParser and AMDGPU.td
The vmem-to-lds-loads-insts feature is only available on gfx9/10.
While target-parser was also enabling it for gfx6,7,8.
---
.../builtins-amdgcn-raw-buffer-load-lds.cl | 12 ++++++++++++
.../CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl | 9 ---------
...ns-amdgcn-raw-ptr-buffer-load-lds-target-error.cl | 4 ++++
llvm/lib/TargetParser/TargetParser.cpp | 2 +-
4 files changed, 17 insertions(+), 10 deletions(-)
create mode 100644
clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-lds.cl
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-lds.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-lds.cl
new file mode 100644
index 0000000000000..8256b61525f9d
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-lds.cl
@@ -0,0 +1,12 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx90a
-emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_lds(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.load.lds(ptr
addrspace(8) [[RSRC:%.*]], ptr addrspace(3) [[LDS:%.*]], i32 1, i32
[[OFFSET:%.*]], i32 [[SOFFSET:%.*]], i32 2, i32 3)
+// CHECK-NEXT: ret void
+//
+void test_amdgcn_raw_ptr_buffer_load_lds(__amdgpu_buffer_rsrc_t rsrc, __local
void * lds, int offset, int soffset) {
+ __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, lds, 1, offset, soffset, 2,
3);
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl
index 5e3ed9027c17a..3403b69e07e4b 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load.cl
@@ -170,12 +170,3 @@ v3u32
test_amdgcn_raw_ptr_buffer_load_b96_non_const_soffset(__amdgpu_buffer_rsrc
v4u32
test_amdgcn_raw_ptr_buffer_load_b128_non_const_soffset(__amdgpu_buffer_rsrc_t
rsrc, int offset, int soffset) {
return __builtin_amdgcn_raw_buffer_load_b128(rsrc, /*offset=*/0, soffset,
/*aux=*/0);
}
-
-// CHECK-LABEL: @test_amdgcn_raw_ptr_buffer_load_lds(
-// CHECK-NEXT: entry:
-// CHECK-NEXT: tail call void @llvm.amdgcn.raw.ptr.buffer.load.lds(ptr
addrspace(8) [[RSRC:%.*]], ptr addrspace(3) [[LDS:%.*]], i32 1, i32
[[OFFSET:%.*]], i32 [[SOFFSET:%.*]], i32 2, i32 3)
-// CHECK-NEXT: ret void
-//
-void test_amdgcn_raw_ptr_buffer_load_lds(__amdgpu_buffer_rsrc_t rsrc, __local
void * lds, int offset, int soffset) {
- __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, lds, 1, offset, soffset, 2,
3);
-}
diff --git
a/clang/test/SemaOpenCL/builtins-amdgcn-raw-ptr-buffer-load-lds-target-error.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-raw-ptr-buffer-load-lds-target-error.cl
index 768f894e9180d..74944f2d93c72 100644
---
a/clang/test/SemaOpenCL/builtins-amdgcn-raw-ptr-buffer-load-lds-target-error.cl
+++
b/clang/test/SemaOpenCL/builtins-amdgcn-raw-ptr-buffer-load-lds-target-error.cl
@@ -1,4 +1,8 @@
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tahiti -S
-verify -o - %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu bonaire -S
-verify -o - %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu carrizo -S
-verify -o - %s
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S
-verify -o - %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -S
-verify -o - %s
// REQUIRES: amdgpu-registered-target
void test_amdgcn_raw_ptr_buffer_load_lds(__amdgpu_buffer_rsrc_t rsrc, __local
void* lds, int offset, int soffset, int x) {
diff --git a/llvm/lib/TargetParser/TargetParser.cpp
b/llvm/lib/TargetParser/TargetParser.cpp
index f7d99ccb57507..7c54901dae47d 100644
--- a/llvm/lib/TargetParser/TargetParser.cpp
+++ b/llvm/lib/TargetParser/TargetParser.cpp
@@ -564,6 +564,7 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const
Triple &T,
case GK_GFX900:
case GK_GFX9_GENERIC:
Features["gfx9-insts"] = true;
+ Features["vmem-to-lds-load-insts"] = true;
[[fallthrough]];
case GK_GFX810:
case GK_GFX805:
@@ -589,7 +590,6 @@ void AMDGPU::fillAMDGPUFeatureMap(StringRef GPU, const
Triple &T,
Features["image-insts"] = true;
Features["s-memtime-inst"] = true;
Features["gws"] = true;
- Features["vmem-to-lds-load-insts"] = true;
break;
case GK_NONE:
break;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
