[clang] [llvm] [AMDGPU] Canonicalize num_records to its actual width in InstCombine (PR #217068)
Krzysztof Drewniak via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 1 11:46:38 PDT 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/217068
>From 10a779c29ed97f04b659d3d46c894139a592fb3c Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Tue, 18 Aug 2026 19:50:17 +0000
Subject: [PATCH] [AMDGPU] Canonicalize num_records to its actual width in
InstCombine
llvm.amdgcn.make.buffer.rsrc is overloaded on the type of its
num_records argument, but the hardware field it ends up in has a fixed
width (32 bits, or 45 bits on gfx1250 and up). Rewrite the intrinsic to
use that width, zero-extending or truncating num_records as needed, so
that IR-level optimizations can see that the extra bits of, for example,
the i64 that Clang emits are not demanded.
Targets that aren't concrete enough for the buffer resource layout to be
known are left alone.
AI disclosure: This was my idea but Claude wrote the code (and I've
tried to tighten up the comments)
---
.../builtins-amdgcn-make-buffer-rsrc.cl | 80 +++++++++----------
.../AMDGPU/AMDGPUInstCombineIntrinsic.cpp | 22 ++++-
.../InstCombine/AMDGPU/amdgcn-intrinsics.ll | 2 +-
.../AMDGPU/make-buffer-rsrc-num-records.ll | 34 +++++---
4 files changed, 80 insertions(+), 58 deletions(-)
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 llvm-commits
mailing list