[llvm] [AMDGPU] Guard intrinsics that need vmem-to-lds-load-insts (PR #214842)
Shilei Tian via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 7 13:00:45 PDT 2026
https://github.com/shiltian created https://github.com/llvm/llvm-project/pull/214842
None
>From 56801baa4cdaa451a6303159ee95e11bb5a78eae Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Fri, 7 Aug 2026 15:53:56 -0400
Subject: [PATCH] [AMDGPU] Guard intrinsics that need vmem-to-lds-load-insts
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 36 ++++++++++++-------
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 2 --
.../AMDGPU/llvm.amdgcn.buffer.load.lds.err.ll | 36 +++++++++++++------
3 files changed, 49 insertions(+), 25 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 9c57b4d343fc6..4ca50ff95eeca 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2071,8 +2071,11 @@ class AMDGPURawBufferLoadLDS : Intrinsic <
// all: volatile op (bit 31, stripped at lowering)
[IntrWillReturn, NoCapture<ArgIndex<1>>, ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<5>>,
ImmArg<ArgIndex<6>>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
-def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS;
-def int_amdgcn_raw_buffer_load_async_lds : AMDGPURawBufferLoadLDS;
+
+let TargetFeatures = "vmem-to-lds-load-insts" in {
+ def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS;
+ def int_amdgcn_raw_buffer_load_async_lds : AMDGPURawBufferLoadLDS;
+} // End TargetFeatures = "vmem-to-lds-load-insts"
class AMDGPURawPtrBufferLoadLDS :
Intrinsic <
@@ -2095,10 +2098,13 @@ class AMDGPURawPtrBufferLoadLDS :
WriteOnly<ArgIndex<1>>, NoCapture<ArgIndex<1>>,
ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<5>>,
ImmArg<ArgIndex<6>>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
-def int_amdgcn_raw_ptr_buffer_load_lds : AMDGPURawPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_lds">;
-def int_amdgcn_raw_ptr_buffer_load_async_lds : AMDGPURawPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_async_lds">;
+
+let TargetFeatures = "vmem-to-lds-load-insts" in {
+ def int_amdgcn_raw_ptr_buffer_load_lds : AMDGPURawPtrBufferLoadLDS,
+ ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_lds">;
+ def int_amdgcn_raw_ptr_buffer_load_async_lds : AMDGPURawPtrBufferLoadLDS,
+ ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_async_lds">;
+} // End TargetFeatures = "vmem-to-lds-load-insts"
class AMDGPUStructBufferLoadLDS : Intrinsic <
@@ -2119,8 +2125,11 @@ class AMDGPUStructBufferLoadLDS : Intrinsic <
// all: volatile op (bit 31, stripped at lowering)
[IntrWillReturn, NoCapture<ArgIndex<1>>, ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<6>>,
ImmArg<ArgIndex<7>>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
-def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS;
-def int_amdgcn_struct_buffer_load_async_lds : AMDGPUStructBufferLoadLDS;
+
+let TargetFeatures = "vmem-to-lds-load-insts" in {
+ def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS;
+ def int_amdgcn_struct_buffer_load_async_lds : AMDGPUStructBufferLoadLDS;
+} // End TargetFeatures = "vmem-to-lds-load-insts"
class AMDGPUStructPtrBufferLoadLDS :
Intrinsic <
@@ -2144,10 +2153,13 @@ class AMDGPUStructPtrBufferLoadLDS :
WriteOnly<ArgIndex<1>>, NoCapture<ArgIndex<1>>,
ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<6>>,
ImmArg<ArgIndex<7>>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
-def int_amdgcn_struct_ptr_buffer_load_lds : AMDGPUStructPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_lds">;
-def int_amdgcn_struct_ptr_buffer_load_async_lds : AMDGPUStructPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_async_lds">;
+
+let TargetFeatures = "vmem-to-lds-load-insts" in {
+ def int_amdgcn_struct_ptr_buffer_load_lds : AMDGPUStructPtrBufferLoadLDS,
+ ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_lds">;
+ def int_amdgcn_struct_ptr_buffer_load_async_lds : AMDGPUStructPtrBufferLoadLDS,
+ ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_async_lds">;
+} // End TargetFeatures = "vmem-to-lds-load-insts"
def int_amdgcn_s_buffer_prefetch_data : DefaultAttrsIntrinsic <
[],
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 1d54903b18def..1d165910adeba 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -12721,8 +12721,6 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
case Intrinsic::amdgcn_struct_buffer_load_async_lds:
case Intrinsic::amdgcn_struct_ptr_buffer_load_lds:
case Intrinsic::amdgcn_struct_ptr_buffer_load_async_lds: {
- if (!Subtarget->hasVMemToLDSLoad())
- return SDValue();
unsigned Opc;
bool HasVIndex =
IntrinsicID == Intrinsic::amdgcn_struct_buffer_load_lds ||
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.load.lds.err.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.load.lds.err.ll
index 44a7a9826163c..dc0373224f0cc 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.load.lds.err.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.load.lds.err.ll
@@ -1,16 +1,30 @@
; RUN: split-file %s %t
;
-; RUN: not --crash llc -mtriple=amdgpu11.00 %t/struct.ll 2>&1 | FileCheck --ignore-case %s
-; RUN: not llc -global-isel -mtriple=amdgpu11.00 %t/struct.ll 2>&1 | FileCheck --ignore-case %s
-; RUN: not --crash llc -mtriple=amdgpu11.00 %t/struct.ptr.ll 2>&1 | FileCheck --ignore-case --check-prefix=LEGALIZER-FAIL %s
-; RUN: not llc -global-isel -mtriple=amdgpu11.00 %t/struct.ptr.ll 2>&1 | FileCheck --ignore-case %s
-; RUN: not --crash llc -mtriple=amdgpu11.00 %t/raw.ll 2>&1 | FileCheck --ignore-case %s
-; RUN: not llc -global-isel -mtriple=amdgpu11.00 %t/raw.ll 2>&1 | FileCheck --ignore-case %s
-; RUN: not --crash llc -mtriple=amdgpu11.00 %t/raw.ptr.ll 2>&1 | FileCheck --ignore-case --check-prefix=LEGALIZER-FAIL %s
-; RUN: not llc -global-isel -mtriple=amdgpu11.00 %t/raw.ptr.ll 2>&1 | FileCheck --ignore-case %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.ll 2>&1 | FileCheck --check-prefix=STRUCT %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.ll 2>&1 | FileCheck --check-prefix=STRUCT %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.async.ll 2>&1 | FileCheck --check-prefix=STRUCT-ASYNC %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.async.ll 2>&1 | FileCheck --check-prefix=STRUCT-ASYNC %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.ptr.ll 2>&1 | FileCheck --check-prefix=STRUCT-PTR %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.ptr.ll 2>&1 | FileCheck --check-prefix=STRUCT-PTR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.ptr.async.ll 2>&1 | FileCheck --check-prefix=STRUCT-PTR-ASYNC %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/struct.ptr.async.ll 2>&1 | FileCheck --check-prefix=STRUCT-PTR-ASYNC %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.ll 2>&1 | FileCheck --check-prefix=RAW %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.ll 2>&1 | FileCheck --check-prefix=RAW %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.async.ll 2>&1 | FileCheck --check-prefix=RAW-ASYNC %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.async.ll 2>&1 | FileCheck --check-prefix=RAW-ASYNC %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.ptr.ll 2>&1 | FileCheck --check-prefix=RAW-PTR %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.ptr.ll 2>&1 | FileCheck --check-prefix=RAW-PTR %s
+; RUN: not llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.ptr.async.ll 2>&1 | FileCheck --check-prefix=RAW-PTR-ASYNC %s
+; RUN: not llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -filetype=null < %t/raw.ptr.async.ll 2>&1 | FileCheck --check-prefix=RAW-PTR-ASYNC %s
;
-; CHECK: LLVM ERROR: Cannot select
-; LEGALIZER-FAIL: Do not know how to expand this operator's operand!
+; STRUCT: llvm.amdgcn.struct.buffer.load.lds requires target feature 'vmem-to-lds-load-insts'
+; STRUCT-ASYNC: llvm.amdgcn.struct.buffer.load.async.lds requires target feature 'vmem-to-lds-load-insts'
+; STRUCT-PTR: llvm.amdgcn.struct.ptr.buffer.load.lds requires target feature 'vmem-to-lds-load-insts'
+; STRUCT-PTR-ASYNC: llvm.amdgcn.struct.ptr.buffer.load.async.lds requires target feature 'vmem-to-lds-load-insts'
+; RAW: llvm.amdgcn.raw.buffer.load.lds requires target feature 'vmem-to-lds-load-insts'
+; RAW-ASYNC: llvm.amdgcn.raw.buffer.load.async.lds requires target feature 'vmem-to-lds-load-insts'
+; RAW-PTR: llvm.amdgcn.raw.ptr.buffer.load.lds requires target feature 'vmem-to-lds-load-insts'
+; RAW-PTR-ASYNC: llvm.amdgcn.raw.ptr.buffer.load.async.lds requires target feature 'vmem-to-lds-load-insts'
;--- struct.ll
define amdgpu_ps void @buffer_load_lds(<4 x i32> inreg %rsrc, ptr addrspace(3) inreg %lds) {
@@ -50,7 +64,7 @@ define amdgpu_ps void @buffer_load_lds(<4 x i32> inreg %rsrc, ptr addrspace(3) i
;--- raw.ptr.ll
define amdgpu_ps void @buffer_load_lds(ptr addrspace(8) inreg %rsrc, ptr addrspace(3) inreg %lds) {
- call void @llvm.amdgcn.raw.ptr.buffer.load.lds(ptr addrspace(8) %rsrc, ptr addrspace(3) %lds, i32 4, i32 0, i32 0, i32 0, i32 0)
+ call void @llvm.amdgcn.raw.ptr.buffer.load.lds(ptr addrspace(8) %rsrc, ptr addrspace(3) %lds, i32 16, i32 0, i32 0, i32 0, i32 0)
ret void
}
More information about the llvm-commits
mailing list