[llvm-branch-commits] [clang] [llvm] [AMDGPU] Canonicalize num_records to its actual width in InstCombine (PR #217068)
Krzysztof Drewniak via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Aug 25 07:57:37 PDT 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/217068
>From f67f89b568ce5a9bccdc69bd1f3839a09125a49a 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 | 54 ++++++-------
.../AMDGPU/make-buffer-rsrc-num-records.ll | 34 +++++---
4 files changed, 106 insertions(+), 84 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 c34a143e4b0b4..5a81475e98711 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
@@ -2177,7 +2177,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 041a67095791c..a5b965fdc48a9 100644
--- a/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll
+++ b/llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll
@@ -43,7 +43,7 @@ define double @test_constant_fold_rcp_f64_1() nounwind {
define float @test_constant_fold_rcp_f32_half() nounwind {
; CHECK-LABEL: @test_constant_fold_rcp_f32_half(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 5.000000e-01) #[[ATTR18:[0-9]+]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 5.000000e-01) #[[ATTR12:[0-9]+]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.rcp.f32(float 0.5) nounwind readnone
@@ -52,7 +52,7 @@ define float @test_constant_fold_rcp_f32_half() nounwind {
define double @test_constant_fold_rcp_f64_half() nounwind {
; CHECK-LABEL: @test_constant_fold_rcp_f64_half(
-; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 5.000000e-01) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 5.000000e-01) #[[ATTR12]]
; CHECK-NEXT: ret double [[VAL]]
;
%val = call double @llvm.amdgcn.rcp.f64(double 0.5) nounwind readnone
@@ -61,7 +61,7 @@ define double @test_constant_fold_rcp_f64_half() nounwind {
define float @test_constant_fold_rcp_f32_43() nounwind {
; CHECK-LABEL: @test_constant_fold_rcp_f32_43(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR12]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) nounwind readnone
@@ -70,7 +70,7 @@ define float @test_constant_fold_rcp_f32_43() nounwind {
define double @test_constant_fold_rcp_f64_43() nounwind {
; CHECK-LABEL: @test_constant_fold_rcp_f64_43(
-; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) #[[ATTR12]]
; CHECK-NEXT: ret double [[VAL]]
;
%val = call double @llvm.amdgcn.rcp.f64(double 4.300000e+01) nounwind readnone
@@ -79,7 +79,7 @@ define double @test_constant_fold_rcp_f64_43() nounwind {
define float @test_constant_fold_rcp_f32_43_strictfp() nounwind strictfp {
; CHECK-LABEL: @test_constant_fold_rcp_f32_43_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR19:[0-9]+]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) #[[ATTR13:[0-9]+]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.rcp.f32(float 4.300000e+01) strictfp nounwind readnone
@@ -190,7 +190,7 @@ define half @test_constant_fold_sqrt_f16_0() nounwind {
define float @test_constant_fold_sqrt_f32_0() nounwind {
; CHECK-LABEL: @test_constant_fold_sqrt_f32_0(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float 0.000000e+00) #[[ATTR12]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.sqrt.f32(float 0.0) nounwind readnone
@@ -199,7 +199,7 @@ define float @test_constant_fold_sqrt_f32_0() nounwind {
define double @test_constant_fold_sqrt_f64_0() nounwind {
; CHECK-LABEL: @test_constant_fold_sqrt_f64_0(
-; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double 0.000000e+00) #[[ATTR12]]
; CHECK-NEXT: ret double [[VAL]]
;
%val = call double @llvm.amdgcn.sqrt.f64(double 0.0) nounwind readnone
@@ -216,7 +216,7 @@ define half @test_constant_fold_sqrt_f16_neg0() nounwind {
define float @test_constant_fold_sqrt_f32_neg0() nounwind {
; CHECK-LABEL: @test_constant_fold_sqrt_f32_neg0(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.sqrt.f32(float -0.000000e+00) #[[ATTR12]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.sqrt.f32(float -0.0) nounwind readnone
@@ -225,7 +225,7 @@ define float @test_constant_fold_sqrt_f32_neg0() nounwind {
define double @test_constant_fold_sqrt_f64_neg0() nounwind {
; CHECK-LABEL: @test_constant_fold_sqrt_f64_neg0(
-; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -0.000000e+00) #[[ATTR18]]
+; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.sqrt.f64(double -0.000000e+00) #[[ATTR12]]
; CHECK-NEXT: ret double [[VAL]]
;
%val = call double @llvm.amdgcn.sqrt.f64(double -0.0) nounwind readnone
@@ -817,7 +817,7 @@ define i1 @test_class_isnan_f32(float %x) nounwind {
define i1 @test_class_isnan_f32_strict(float %x) nounwind strictfp {
; CHECK-LABEL: @test_class_isnan_f32_strict(
-; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (nan) */ i32 3) #[[ATTR20:[0-9]+]]
+; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (nan) */ i32 3) #[[ATTR10:[0-9]+]]
; CHECK-NEXT: ret i1 [[VAL]]
;
%val = call i1 @llvm.amdgcn.class.f32(float %x, i32 3) strictfp
@@ -835,7 +835,7 @@ define i1 @test_class_is_p0_n0_f32(float %x) nounwind {
define i1 @test_class_is_p0_n0_f32_strict(float %x) nounwind strictfp {
; CHECK-LABEL: @test_class_is_p0_n0_f32_strict(
-; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (zero) */ i32 96) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call i1 @llvm.is.fpclass.f32(float [[X:%.*]], /* (zero) */ i32 96) #[[ATTR10]]
; CHECK-NEXT: ret i1 [[VAL]]
;
%val = call i1 @llvm.amdgcn.class.f32(float %x, i32 96) strictfp
@@ -5843,7 +5843,7 @@ define double @trig_preop_constfold_neg32_segment() {
define double @trig_preop_constfold_strictfp() strictfp {
; CHECK-LABEL: @trig_preop_constfold_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double 3.454350e+02, i32 5) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double 3.454350e+02, i32 5) #[[ATTR10]]
; CHECK-NEXT: ret double [[VAL]]
;
%val = call double @llvm.amdgcn.trig.preop.f64(double 3.454350e+02, i32 5) strictfp
@@ -6137,7 +6137,7 @@ define double @trig_preop_strip_copysign(double %mag, double %sign, i32 %idx) {
define double @trig_preop_strip_fabs_strictfp(double %val, i32 %idx) strictfp {
; CHECK-LABEL: @trig_preop_strip_fabs_strictfp(
-; CHECK-NEXT: [[RESULT:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double [[VAL:%.*]], i32 [[IDX:%.*]]) #[[ATTR20]]
+; CHECK-NEXT: [[RESULT:%.*]] = call double @llvm.amdgcn.trig.preop.f64(double [[VAL:%.*]], i32 [[IDX:%.*]]) #[[ATTR10]]
; CHECK-NEXT: ret double [[RESULT]]
;
%fabs = call double @llvm.fabs.f64(double %val)
@@ -6276,7 +6276,7 @@ define half @test_constant_fold_log_f16_neg10() {
define float @test_constant_fold_log_f32_qnan_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_log_f32_qnan_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float +qnan) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float +qnan) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.log.f32(float 0x7FF8000000000000) strictfp
@@ -6285,7 +6285,7 @@ define float @test_constant_fold_log_f32_qnan_strictfp() strictfp {
define float @test_constant_fold_log_f32_0_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_log_f32_0_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float 0.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float 0.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.log.f32(float 0.0) strictfp
@@ -6294,7 +6294,7 @@ define float @test_constant_fold_log_f32_0_strictfp() strictfp {
define float @test_constant_fold_log_f32_neg0_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_log_f32_neg0_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -0.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -0.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.log.f32(float -0.0) strictfp
@@ -6303,7 +6303,7 @@ define float @test_constant_fold_log_f32_neg0_strictfp() strictfp {
define float @test_constant_fold_log_f32_neg_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_log_f32_neg_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -1.000000e+01) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -1.000000e+01) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.log.f32(float -10.0) strictfp
@@ -6320,7 +6320,7 @@ define float @test_constant_fold_log_f32_pinf_strictfp() strictfp {
define float @test_constant_fold_log_f32_ninf_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_log_f32_ninf_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -inf) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.log.f32(float -inf) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.log.f32(float 0xFFF0000000000000) strictfp
@@ -6522,7 +6522,7 @@ define half @test_constant_fold_exp2_f16_neg10() {
define float @test_constant_fold_exp2_f32_qnan_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_qnan_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float +qnan) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float +qnan) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float 0x7FF8000000000000) strictfp
@@ -6531,7 +6531,7 @@ define float @test_constant_fold_exp2_f32_qnan_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_0_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_0_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 0.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 0.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float 0.0) strictfp
@@ -6540,7 +6540,7 @@ define float @test_constant_fold_exp2_f32_0_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_neg0_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_neg0_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -0.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -0.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float -0.0) strictfp
@@ -6549,7 +6549,7 @@ define float @test_constant_fold_exp2_f32_neg0_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_1_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_1_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 1.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 1.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float 1.0) strictfp
@@ -6558,7 +6558,7 @@ define float @test_constant_fold_exp2_f32_1_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_neg1_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_neg1_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float -1.0) strictfp
@@ -6567,7 +6567,7 @@ define float @test_constant_fold_exp2_f32_neg1_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_2_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_2_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 2.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float 2.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float 2.0) strictfp
@@ -6576,7 +6576,7 @@ define float @test_constant_fold_exp2_f32_2_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_neg2_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_neg2_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -2.000000e+00) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -2.000000e+00) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float -2.0) strictfp
@@ -6585,7 +6585,7 @@ define float @test_constant_fold_exp2_f32_neg2_strictfp() strictfp {
define float @test_constant_fold_exp2_f32_neg_strictfp() strictfp {
; CHECK-LABEL: @test_constant_fold_exp2_f32_neg_strictfp(
-; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+01) #[[ATTR20]]
+; CHECK-NEXT: [[VAL:%.*]] = call float @llvm.amdgcn.exp2.f32(float -1.000000e+01) #[[ATTR10]]
; CHECK-NEXT: ret float [[VAL]]
;
%val = call float @llvm.amdgcn.exp2.f32(float -10.0) strictfp
@@ -6699,7 +6699,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-branch-commits
mailing list