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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
