[clang] 23bc218 - [AMDGPU] Canonicalize num_records to its actual width in InstCombine (#217068)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 1 13:47:14 PDT 2026
Author: Krzysztof Drewniak
Date: 2026-10-01T13:47:05-07:00
New Revision: 23bc2182f84a42e3bc108d5115bd272faed545c5
URL: https://github.com/llvm/llvm-project/commit/23bc2182f84a42e3bc108d5115bd272faed545c5
DIFF: https://github.com/llvm/llvm-project/commit/23bc2182f84a42e3bc108d5115bd272faed545c5.diff
LOG: [AMDGPU] Canonicalize num_records to its actual width in InstCombine (#217068)
This PR adds code to InstCombineIntrinsic to change the width of the
num_recods field (by extension or truncation) to the correct width for
the target triple (if a concrete enough target triple has been set) so
that LLVM IR-level optimizations can see the lack of demand for the
high bits, for example.
Assisted by Claude, which also found those buffer lowering edge cases
Added:
Modified:
clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll
llvm/test/Transforms/InstCombine/AMDGPU/make-buffer-rsrc-num-records.ll
Removed:
################################################################################
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
index 5ebc4f68f2922..d794f2c746dee 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
@@ -6,16 +6,15 @@
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0(
// GFX601-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[NUM]], i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0(
// GFX1250-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[TMP0]], i32 [[FLAGS]])
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p, short stride, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
@@ -24,16 +23,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p, short stride, in
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0_stride_constant(
// GFX601-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 4, i64 [[CONV]], i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 4, i32 [[NUM]], i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0_stride_constant(
// GFX1250-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 4, i64 [[CONV]], i32 [[FLAGS]])
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 4, i45 [[TMP0]], i32 [[FLAGS]])
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_stride_constant(void *p, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
@@ -42,13 +40,13 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_stride_constant(void *p,
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0_num_constant(
// GFX601-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 1234, i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 1234, i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0_num_constant(
// GFX1250-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 1234, i32 [[FLAGS]])
+// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 1234, i32 [[FLAGS]])
// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p, short stride, int flags) {
@@ -58,16 +56,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p, sho
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0_flags_constant(
// GFX601-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 5678)
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[NUM]], i32 5678)
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p0_flags_constant(
// GFX1250-SAME: ptr nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 5678)
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[TMP0]], i32 5678)
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p, short stride, int num) {
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
@@ -76,16 +73,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p, s
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1(
// GFX601-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i32(ptr addrspace(1) [[P]], i16 [[STRIDE]], i32 [[NUM]], i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1(
// GFX1250-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i45(ptr addrspace(1) [[P]], i16 [[STRIDE]], i45 [[TMP0]], i32 [[FLAGS]])
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1(global void *p, short stride, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
@@ -94,16 +90,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1(global void *p, short str
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1_stride_constant(
// GFX601-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 4, i64 [[CONV]], i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i32(ptr addrspace(1) [[P]], i16 4, i32 [[NUM]], i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1_stride_constant(
// GFX1250-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 4, i64 [[CONV]], i32 [[FLAGS]])
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i45(ptr addrspace(1) [[P]], i16 4, i45 [[TMP0]], i32 [[FLAGS]])
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_stride_constant(global void *p, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
@@ -112,13 +107,13 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_stride_constant(global vo
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1_num_constant(
// GFX601-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 [[STRIDE]], i64 1234, i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i32(ptr addrspace(1) [[P]], i16 [[STRIDE]], i32 1234, i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1_num_constant(
// GFX1250-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 [[STRIDE]], i64 1234, i32 [[FLAGS]])
+// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i45(ptr addrspace(1) [[P]], i16 [[STRIDE]], i45 1234, i32 [[FLAGS]])
// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_num_constant(global void *p, short stride, int flags) {
@@ -128,16 +123,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_num_constant(global void
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1_flags_constant(
// GFX601-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 5678)
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i32(ptr addrspace(1) [[P]], i16 [[STRIDE]], i32 [[NUM]], i32 5678)
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_rsrc_p1_flags_constant(
// GFX1250-SAME: ptr addrspace(1) nofree noundef readnone captures(ret: address, provenance) [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) [[P]], i16 [[STRIDE]], i64 [[CONV]], i32 5678)
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i45(ptr addrspace(1) [[P]], i16 [[STRIDE]], i45 [[TMP0]], i32 5678)
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_flags_constant(global void *p, short stride, int num) {
return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
@@ -146,16 +140,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_flags_constant(global voi
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_p0_nullptr(
// GFX601-SAME: i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr null, i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr null, i16 [[STRIDE]], i32 [[NUM]], i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_p0_nullptr(
// GFX1250-SAME: i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr null, i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr null, i16 [[STRIDE]], i45 [[TMP0]], i32 [[FLAGS]])
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p0_nullptr(short stride, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc((void *)0LL, stride, num, flags);
@@ -164,16 +157,15 @@ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p0_nullptr(short stride, int num,
// GFX601-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_p1_nullptr(
// GFX601-SAME: i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX601-NEXT: [[ENTRY:.*:]]
-// GFX601-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) null, i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
+// GFX601-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i32(ptr addrspace(1) null, i16 [[STRIDE]], i32 [[NUM]], i32 [[FLAGS]])
// GFX601-NEXT: ret ptr addrspace(8) [[TMP0]]
//
// GFX1250-LABEL: define dso_local noundef ptr addrspace(8) @test_amdgcn_make_buffer_p1_nullptr(
// GFX1250-SAME: i16 noundef signext [[STRIDE:%.*]], i32 noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) local_unnamed_addr #[[ATTR0]] {
// GFX1250-NEXT: [[ENTRY:.*:]]
-// GFX1250-NEXT: [[CONV:%.*]] = sext i32 [[NUM]] to i64
-// GFX1250-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) null, i16 [[STRIDE]], i64 [[CONV]], i32 [[FLAGS]])
-// GFX1250-NEXT: ret ptr addrspace(8) [[TMP0]]
+// GFX1250-NEXT: [[TMP0:%.*]] = sext i32 [[NUM]] to i45
+// GFX1250-NEXT: [[TMP1:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i45(ptr addrspace(1) null, i16 [[STRIDE]], i45 [[TMP0]], i32 [[FLAGS]])
+// GFX1250-NEXT: ret ptr addrspace(8) [[TMP1]]
//
__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p1_nullptr(short stride, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc((global void *)0LL, stride, num, flags);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
index 4936b16148ce4..6b5976bab177a 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
@@ -2132,7 +2132,27 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
Value *Src = II.getArgOperand(0);
if (isa<PoisonValue>(Src))
return IC.replaceInstUsesWith(II, PoisonValue::get(II.getType()));
- return std::nullopt;
+
+ // Normalize num_records to the correct width.
+ std::optional<unsigned> Width = ST->getBufferResourceNumRecordsWidth();
+ if (!Width)
+ return std::nullopt;
+ Type *NumRecordsTy = IC.Builder.getIntNTy(*Width);
+ if (II.getArgOperand(2)->getType() == NumRecordsTy)
+ return std::nullopt;
+ SmallVector<Value *, 4> Args(II.args());
+ Args[2] = IC.Builder.CreateZExtOrTrunc(Args[2], NumRecordsTy);
+ CallInst *NewCall = IC.Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::amdgcn_make_buffer_rsrc,
+ {II.getType(), Src->getType(), NumRecordsTy}, Args);
+ NewCall->copyMetadata(II);
+ NewCall->setTailCallKind(II.getTailCallKind());
+ // Copy over all attributes except those on num_records, which may no longer
+ // be valid.
+ NewCall->setAttributes(
+ II.getAttributes().removeParamAttributes(II.getContext(), 2));
+ NewCall->takeName(&II);
+ return IC.replaceInstUsesWith(II, NewCall);
}
case Intrinsic::amdgcn_raw_buffer_store_format:
case Intrinsic::amdgcn_struct_buffer_store_format:
diff --git a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll
index 9cd0ee45bdfc1..fb00d8e0ad3a8 100644
--- a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll
+++ b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll
@@ -5960,7 +5960,7 @@ define ptr addrspace(8) @make_buffer_rsrc_poison() {
define ptr addrspace(8) @make_buffer_rsrc_undef() {
; CHECK-LABEL: @make_buffer_rsrc_undef(
-; CHECK-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) undef, i16 0, i64 1234, i32 5678)
+; CHECK-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i32(ptr addrspace(1) undef, i16 0, i32 1234, i32 5678)
; CHECK-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p1.i64(ptr addrspace(1) undef, i16 0, i64 1234, i32 5678)
diff --git a/llvm/test/Transforms/InstCombine/AMDGPU/make-buffer-rsrc-num-records.ll b/llvm/test/Transforms/InstCombine/AMDGPU/make-buffer-rsrc-num-records.ll
index 7a6e6b43e345e..1e373e8cf7711 100644
--- a/llvm/test/Transforms/InstCombine/AMDGPU/make-buffer-rsrc-num-records.ll
+++ b/llvm/test/Transforms/InstCombine/AMDGPU/make-buffer-rsrc-num-records.ll
@@ -11,12 +11,14 @@ define ptr addrspace(8) @from_i64(ptr %p, i16 %stride, i64 %num, i32 %flags) {
;
; RECORDS32-LABEL: define ptr addrspace(8) @from_i64(
; RECORDS32-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 [[NUM]], i32 [[FLAGS]])
+; RECORDS32-NEXT: [[TMP1:%.*]] = trunc i64 [[NUM]] to i32
+; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[TMP1]], i32 [[FLAGS]])
; RECORDS32-NEXT: ret ptr addrspace(8) [[RSRC]]
;
; RECORDS45-LABEL: define ptr addrspace(8) @from_i64(
; RECORDS45-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 [[NUM]], i32 [[FLAGS]])
+; RECORDS45-NEXT: [[TMP1:%.*]] = trunc i64 [[NUM]] to i45
+; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[TMP1]], i32 [[FLAGS]])
; RECORDS45-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr %p, i16 %stride, i64 %num, i32 %flags)
@@ -31,12 +33,14 @@ define ptr addrspace(8) @from_i128(ptr %p, i16 %stride, i128 %num, i32 %flags) {
;
; RECORDS32-LABEL: define ptr addrspace(8) @from_i128(
; RECORDS32-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i128 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i128(ptr [[P]], i16 [[STRIDE]], i128 [[NUM]], i32 [[FLAGS]])
+; RECORDS32-NEXT: [[TMP1:%.*]] = trunc i128 [[NUM]] to i32
+; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[TMP1]], i32 [[FLAGS]])
; RECORDS32-NEXT: ret ptr addrspace(8) [[RSRC]]
;
; RECORDS45-LABEL: define ptr addrspace(8) @from_i128(
; RECORDS45-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i128 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i128(ptr [[P]], i16 [[STRIDE]], i128 [[NUM]], i32 [[FLAGS]])
+; RECORDS45-NEXT: [[TMP1:%.*]] = trunc i128 [[NUM]] to i45
+; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[TMP1]], i32 [[FLAGS]])
; RECORDS45-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i128(ptr %p, i16 %stride, i128 %num, i32 %flags)
@@ -56,7 +60,8 @@ define ptr addrspace(8) @from_i32(ptr %p, i16 %stride, i32 %num, i32 %flags) {
;
; RECORDS45-LABEL: define ptr addrspace(8) @from_i32(
; RECORDS45-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[NUM]], i32 [[FLAGS]])
+; RECORDS45-NEXT: [[TMP1:%.*]] = zext i32 [[NUM]] to i45
+; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[TMP1]], i32 [[FLAGS]])
; RECORDS45-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr %p, i16 %stride, i32 %num, i32 %flags)
@@ -71,7 +76,8 @@ define ptr addrspace(8) @from_i45(ptr %p, i16 %stride, i45 %num, i32 %flags) {
;
; RECORDS32-LABEL: define ptr addrspace(8) @from_i45(
; RECORDS32-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i45 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[NUM]], i32 [[FLAGS]])
+; RECORDS32-NEXT: [[TMP1:%.*]] = trunc i45 [[NUM]] to i32
+; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[TMP1]], i32 [[FLAGS]])
; RECORDS32-NEXT: ret ptr addrspace(8) [[RSRC]]
;
; RECORDS45-LABEL: define ptr addrspace(8) @from_i45(
@@ -91,12 +97,14 @@ define ptr addrspace(8) @from_i16(ptr %p, i16 %stride, i16 %num, i32 %flags) {
;
; RECORDS32-LABEL: define ptr addrspace(8) @from_i16(
; RECORDS32-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i16 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i16(ptr [[P]], i16 [[STRIDE]], i16 [[NUM]], i32 [[FLAGS]])
+; RECORDS32-NEXT: [[TMP1:%.*]] = zext i16 [[NUM]] to i32
+; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 [[TMP1]], i32 [[FLAGS]])
; RECORDS32-NEXT: ret ptr addrspace(8) [[RSRC]]
;
; RECORDS45-LABEL: define ptr addrspace(8) @from_i16(
; RECORDS45-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i16 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i16(ptr [[P]], i16 [[STRIDE]], i16 [[NUM]], i32 [[FLAGS]])
+; RECORDS45-NEXT: [[TMP1:%.*]] = zext i16 [[NUM]] to i45
+; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 [[TMP1]], i32 [[FLAGS]])
; RECORDS45-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i16(ptr %p, i16 %stride, i16 %num, i32 %flags)
@@ -111,12 +119,12 @@ define ptr addrspace(8) @from_constant(ptr %p, i16 %stride, i32 %flags) {
;
; RECORDS32-LABEL: define ptr addrspace(8) @from_constant(
; RECORDS32-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 4294967296, i32 [[FLAGS]])
+; RECORDS32-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr [[P]], i16 [[STRIDE]], i32 0, i32 [[FLAGS]])
; RECORDS32-NEXT: ret ptr addrspace(8) [[RSRC]]
;
; RECORDS45-LABEL: define ptr addrspace(8) @from_constant(
; RECORDS45-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr [[P]], i16 [[STRIDE]], i64 4294967296, i32 [[FLAGS]])
+; RECORDS45-NEXT: [[RSRC:%.*]] = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr [[P]], i16 [[STRIDE]], i45 4294967296, i32 [[FLAGS]])
; RECORDS45-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr %p, i16 %stride, i64 4294967296, i32 %flags)
@@ -131,12 +139,14 @@ define ptr addrspace(8) @keeps_metadata_and_attributes(ptr %p, i16 %stride, i64
;
; RECORDS32-LABEL: define ptr addrspace(8) @keeps_metadata_and_attributes(
; RECORDS32-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS32-NEXT: [[RSRC:%.*]] = tail call align 16 ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr noundef [[P]], i16 [[STRIDE]], i64 [[NUM]], i32 [[FLAGS]]), !amdgpu.uniform [[META0:![0-9]+]]
+; RECORDS32-NEXT: [[TMP1:%.*]] = trunc i64 [[NUM]] to i32
+; RECORDS32-NEXT: [[RSRC:%.*]] = tail call align 16 ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i32(ptr noundef [[P]], i16 [[STRIDE]], i32 [[TMP1]], i32 [[FLAGS]]), !amdgpu.uniform [[META0:![0-9]+]]
; RECORDS32-NEXT: ret ptr addrspace(8) [[RSRC]]
;
; RECORDS45-LABEL: define ptr addrspace(8) @keeps_metadata_and_attributes(
; RECORDS45-SAME: ptr [[P:%.*]], i16 [[STRIDE:%.*]], i64 [[NUM:%.*]], i32 [[FLAGS:%.*]]) {
-; RECORDS45-NEXT: [[RSRC:%.*]] = tail call align 16 ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr noundef [[P]], i16 [[STRIDE]], i64 [[NUM]], i32 [[FLAGS]]), !amdgpu.uniform [[META0:![0-9]+]]
+; RECORDS45-NEXT: [[TMP1:%.*]] = trunc i64 [[NUM]] to i45
+; RECORDS45-NEXT: [[RSRC:%.*]] = tail call align 16 ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i45(ptr noundef [[P]], i16 [[STRIDE]], i45 [[TMP1]], i32 [[FLAGS]]), !amdgpu.uniform [[META0:![0-9]+]]
; RECORDS45-NEXT: ret ptr addrspace(8) [[RSRC]]
;
%rsrc = tail call align 16 ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p8.p0.i64(ptr noundef %p, i16 %stride, i64 %num, i32 %flags), !amdgpu.uniform !0
More information about the cfe-commits
mailing list