Skip to content

[AMDGPU] Add ptr.s.buffer.load intrinsic, use it from Clang #209243

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
krzysz00 wants to merge 2 commits into
base: main
Choose a base branch
from
Open

[AMDGPU] Add ptr.s.buffer.load intrinsic, use it from Clang #209243

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0146f26
[AMDGPU] Add ptr.s.buffer.load intrinsic, use it from Clang
krzysz00 Jul 10, 2026
eda7aa7
Merge branch 'main' into ptr-s-buffer-load
krzysz00 Jul 13, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 2 additions & 3 deletions clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -718,7 +718,7 @@ value gives an undefined result.

def DocCatSBufferLoad : DocumentationCategory<"S-Buffer Load Builtins"> {
let Content = [{
These builtins lower to ``llvm.amdgcn.s.buffer.load`` and issue an
These builtins lower to ``llvm.amdgcn.ptr.s.buffer.load`` and issue an
``s_buffer_load`` when the byte offset is uniform across the wavefront.
When the offset is divergent, the backend may lower to a ``buffer_load``.

Expand All @@ -735,8 +735,7 @@ When the offset is divergent, the backend may lower to a ``buffer_load``.

These builtins take the buffer resource as ``__amdgpu_buffer_rsrc_t``, which
can be created with ``__builtin_amdgcn_make_buffer_rsrc``. The backend
lowers the resource to the 4-dword SGPR descriptor required by
``llvm.amdgcn.s.buffer.load``.
lowers the resource to the 4-dword SGPR descriptor required by the instruction.

The return type selects the load width. Separate builtins are provided for
each supported scalar and vector element type.
Expand Down
17 changes: 8 additions & 9 deletions clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,17 +34,16 @@ namespace {
static Value *emitAMDGPUSBufferLoadBuiltin(CodeGenFunction &CGF,
const CallExpr *E) {
llvm::Type *RetTy = CGF.ConvertType(E->getType());
Function *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_s_buffer_load, RetTy);
Function *F =
CGF.CGM.getIntrinsic(Intrinsic::amdgcn_ptr_s_buffer_load, RetTy);

Value *RsrcPtr = CGF.EmitScalarExpr(E->getArg(0));
llvm::Type *I128Ty = llvm::IntegerType::get(CGF.getLLVMContext(), 128);
llvm::Type *RsrcVecTy =
llvm::FixedVectorType::get(CGF.Builder.getInt32Ty(), 4);
Value *RsrcInt = CGF.Builder.CreatePtrToInt(RsrcPtr, I128Ty);
Value *Rsrc = CGF.Builder.CreateBitCast(RsrcInt, RsrcVecTy);

return CGF.Builder.CreateCall(F, {Rsrc, CGF.EmitScalarExpr(E->getArg(1)),
CGF.EmitScalarExpr(E->getArg(2))});
CallInst *Call =
CGF.Builder.CreateCall(F, {RsrcPtr, CGF.EmitScalarExpr(E->getArg(1)),
CGF.EmitScalarExpr(E->getArg(2))});
Call->setMetadata(llvm::LLVMContext::MD_invariant_load,
llvm::MDNode::get(CGF.Builder.getContext(), {}));
return Call;
}

// Has second type mangled argument.
Expand Down
96 changes: 24 additions & 72 deletions clang/test/CodeGenOpenCL/builtins-amdgcn-s-buffer-load.cl
View file
Open in desktop

Large diffs are not rendered by default.

27 changes: 27 additions & 0 deletions llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1485,6 +1485,33 @@ def int_amdgcn_s_buffer_load : DefaultAttrsIntrinsic <
[IntrNoMem, ImmArg<ArgIndex<2>>]>,
AMDGPURsrcIntrinsic<0>;

// Generate an s_buffer_load instruction, falling back to buffer_load if the offset
// is not uniform. This operates like the s.buffer.load intrinsic, but uses a buffer

@shiltian shiltian Jul 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why does int_amdgcn_s_buffer_load do the other way around, aka generate buffer_load by default but optimize to s_buffer_load?

@arsenm arsenm Jul 13, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

the scalar cache isn't coherent with vector, it's difficult to migrate that direction

@krzysz00 krzysz00 Jul 13, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think s_buffer_load is documented weird, given that it's got those "no memory" semantics and all.

Both of these are an s_buffer_load when all the arguments are uniform and a buffer_load otherwise, I just wrote the docs on this one as "you're explicitly requesting a uniform load".

@krzysz00 krzysz00 Jul 13, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

the scalar cache isn't coherent with vector, it's difficult to migrate that direction

Which is why (though this was underdocumented for the "old form") you must only s_buffer_load invariant things/constants, therefore preventing the coherence issue from coming up in any realistic workload

@slinder1 slinder1 Jul 16, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should probably update the old intrinsic wording too, then?

Alternatively it seems like we can define one in terms of the other. I.e. keep your now wording, and update the comment for the old intrinsic to be "an int_amdgcn_ptr_s_buffer_load except (describe the difference)..."

@krzysz00 krzysz00 Jul 16, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah, I'll swap the old wording to point to the new one

// resource pointer (`ptr addrspace(8)`) to hold the resource (which should be

@slinder1 slinder1 Jul 16, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is the "should" here a requirement, or a suggestion?

// wave-uniform) instead of a `<4 x i32>`. Note that some architectures require
// the offset to be aligned to the width of the type being loaded (or, in some cases
// the shorter of that alignment and 4 bytes).
//
// Also note that this data is loaded via a distinct cache than typical load
// instructions or other buffer_load operations. Therefore, the data pointed to by
// this load *must* be invariant to avoid coherence risks. Frontends emitting
// this intrinsic *should* mark it `!invariant.load`, and the backend may also
// add such markings.
def int_amdgcn_ptr_s_buffer_load : DefaultAttrsIntrinsic <
[llvm_any_ty],
[AMDGPUBufferRsrcTy, // rsrc(SGPR)
llvm_i32_ty, // byte offset
llvm_i32_ty], // auxiliary/cachepolicy(imm):
// bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
// bit 3 = swz, bit 4 = scc (gfx90a)
// gfx942: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
// gfx12+: bits [0-2] = th, bits [3-4] = scope,
// bit 6 = swz
// Note: volatile bit is **not** permitted here.
[IntrArgMemOnly, IntrReadMem, ReadOnly<ArgIndex<0>>,
NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<2>>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;

// Buffer intrinsics with separate raw and struct variants. The raw
// variant never has an index. The struct variant always has an index, even if
// it is const 0. A struct intrinsic with constant 0 index is different to the
Expand Down
1 change: 1 addition & 0 deletions llvm/lib/Target/AMDGPU/AMDGPUAsanInstrumentation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -323,6 +323,7 @@ void getInterestingMemoryOperands(
case Intrinsic::amdgcn_struct_tbuffer_load:
case Intrinsic::amdgcn_struct_ptr_tbuffer_load:
case Intrinsic::amdgcn_s_buffer_load:
case Intrinsic::amdgcn_ptr_s_buffer_load:
case Intrinsic::amdgcn_global_load_tr_b64:
case Intrinsic::amdgcn_global_load_tr_b128: {
unsigned PtrOpNo = 0;
Expand Down
2 changes: 2 additions & 0 deletions llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2328,6 +2328,7 @@ static Value *simplifyAMDGCNMemoryIntrinsicDemanded(InstCombiner &IC,
OffsetIdx = 1;
break;
case Intrinsic::amdgcn_s_buffer_load:
case Intrinsic::amdgcn_ptr_s_buffer_load:
// If resulting type is vec3, there is no point in trimming the
// load with updated offset, as the vec3 would most likely be widened to
// vec4 anyway during lowering.
Expand Down Expand Up @@ -2530,6 +2531,7 @@ std::optional<Value *> GCNTTIImpl::simplifyDemandedVectorEltsIntrinsic(
case Intrinsic::amdgcn_raw_tbuffer_load:
case Intrinsic::amdgcn_raw_ptr_tbuffer_load:
case Intrinsic::amdgcn_s_buffer_load:
case Intrinsic::amdgcn_ptr_s_buffer_load:
case Intrinsic::amdgcn_struct_buffer_load:
case Intrinsic::amdgcn_struct_ptr_buffer_load:
case Intrinsic::amdgcn_struct_buffer_load_format:
Expand Down
33 changes: 18 additions & 15 deletions llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7665,6 +7665,7 @@ bool AMDGPULegalizerInfo::legalizeSBufferLoad(LegalizerHelper &Helper,
LLT Ty = B.getMRI()->getType(OrigDst);
unsigned Size = Ty.getSizeInBits();
MachineFunction &MF = B.getMF();
bool HasMMO = !MI.memoperands_empty();
unsigned Opc = 0;
if (Size < 32 && ST.hasScalarSubwordLoads()) {
assert(Size == 8 || Size == 16);
Expand All @@ -7691,22 +7692,23 @@ bool AMDGPULegalizerInfo::legalizeSBufferLoad(LegalizerHelper &Helper,
B.setInsertPt(B.getMBB(), MI);
}

// FIXME: We don't really need this intermediate instruction. The intrinsic
// should be fixed to have a memory operand. Since it's readnone, we're not
// allowed to add one.
MI.setDesc(B.getTII().get(Opc));
MI.removeOperand(1); // Remove intrinsic ID

// FIXME: When intrinsic definition is fixed, this should have an MMO already.
const unsigned MemSize = (Size + 7) / 8;
const Align MemAlign = B.getDataLayout().getABITypeAlign(
getTypeForLLT(Ty, MF.getFunction().getContext()));
MachineMemOperand *MMO = MF.getMachineMemOperand(
MachinePointerInfo(),
MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
MachineMemOperand::MOInvariant,
MemSize, MemAlign);
MI.addMemOperand(MF, MMO);
MI.removeOperand(1);
castBufferRsrcArgToV4I32(MI, B, 1);

if (!HasMMO) {
// Legacy intrinsic that doesn't take a pointer and so can't already have an
// MMO.
const unsigned MemSize = (Size + 7) / 8;
const Align MemAlign = B.getDataLayout().getABITypeAlign(
getTypeForLLT(Ty, MF.getFunction().getContext()));
MachineMemOperand *MMO = MF.getMachineMemOperand(
MachinePointerInfo(),
MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
MachineMemOperand::MOInvariant,
MemSize, MemAlign);
MI.addMemOperand(MF, MMO);
}
if (Dst != OrigDst) {
MI.getOperand(0).setReg(Dst);
B.setInsertPt(B.getMBB(), ++B.getInsertPt());
Expand Down Expand Up @@ -8492,6 +8494,7 @@ bool AMDGPULegalizerInfo::legalizeIntrinsic(LegalizerHelper &Helper,
return true;
}
case Intrinsic::amdgcn_s_buffer_load:
case Intrinsic::amdgcn_ptr_s_buffer_load:
return legalizeSBufferLoad(Helper, MI);
case Intrinsic::amdgcn_raw_buffer_store:
case Intrinsic::amdgcn_raw_ptr_buffer_store:
Expand Down
16 changes: 16 additions & 0 deletions llvm/lib/Target/AMDGPU/AMDGPULowerIntrinsics.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,7 @@ class AMDGPULowerIntrinsicsImpl {

private:
bool visitBarrier(IntrinsicInst &I);
bool visitPtrSBufferLoad(IntrinsicInst &I);
};

class AMDGPULowerIntrinsicsLegacy : public ModulePass {
Expand Down Expand Up @@ -76,6 +77,10 @@ bool AMDGPULowerIntrinsicsImpl::run() {
case Intrinsic::amdgcn_s_cluster_barrier:
forEachCall(F, [&](IntrinsicInst *II) { Changed |= visitBarrier(*II); });
break;
case Intrinsic::amdgcn_ptr_s_buffer_load:
forEachCall(
F, [&](IntrinsicInst *II) { Changed |= visitPtrSBufferLoad(*II); });
break;
}
}

Expand Down Expand Up @@ -193,6 +198,17 @@ bool AMDGPULowerIntrinsicsImpl::visitBarrier(IntrinsicInst &I) {
return false;
}

bool AMDGPULowerIntrinsicsImpl::visitPtrSBufferLoad(IntrinsicInst &I) {
assert(I.getIntrinsicID() == Intrinsic::amdgcn_ptr_s_buffer_load);

if (I.hasMetadata(LLVMContext::MD_invariant_load))
return false;

I.setMetadata(LLVMContext::MD_invariant_load,
MDNode::get(I.getContext(), {}));
return true;
}

PreservedAnalyses AMDGPULowerIntrinsicsPass::run(Module &M,
ModuleAnalysisManager &MAM) {
AMDGPULowerIntrinsicsImpl Impl(M, TM);
Expand Down
18 changes: 8 additions & 10 deletions llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -644,19 +644,16 @@ bool RegBankLegalizeHelper::lowerSBufToBuf(MachineInstr &MI,
NumLoads = LoadSize / 128;
Ty = Ty.divide(NumLoads);
}
for (int i = 0; i < NumLoads; ++i)
LoadParts.emplace_back(MRI.createVirtualRegister({VgprRB, Ty}));
MachineMemOperand *OrigMMO = *MI.memoperands_begin();
for (int I = 0; I < NumLoads; ++I)

@slinder1 slinder1 Jul 16, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there a side-effect here that makes moving this snippet meaningful?

LoadParts.emplace_back(MRI.createVirtualRegister({VgprRB, Ty}));
const Align Alignment = OrigMMO->getAlign();
MachineFunction &MF = B.getMF();
Register SOffset;
Register VOffset;
int64_t ImmOffset = 0;
unsigned MMOOffset = setBufferOffsets(B, MI.getOperand(2).getReg(), VOffset,
SOffset, ImmOffset, Alignment);
// Use the MMO size from the original instruction rather than the (possibly
// widened) register type. E.g. 96-bit loads are widened to 128-bit during
// legalization but the MMO still reflects the original 96-bit access size.
const unsigned MemSize = divideCeil(OrigMMO->getSize().getValue(), NumLoads);
MachineMemOperand *BaseMMO = MF.getMachineMemOperand(OrigMMO, 0, MemSize);
if (MMOOffset != 0)
Expand All @@ -665,6 +662,7 @@ bool RegBankLegalizeHelper::lowerSBufToBuf(MachineInstr &MI,
// instead. We can assume that the buffer is unswizzled.
Register RSrc = MI.getOperand(1).getReg();
Register VIndex = B.buildConstant(VgprRB_I32, 0).getReg(0);
unsigned CachePolicy = MI.getOperand(3).getImm();
unsigned Opc = AMDGPU::G_AMDGPU_BUFFER_LOAD;
switch (MI.getOpcode()) {
case AMDGPU::G_AMDGPU_S_BUFFER_LOAD_SBYTE:
Expand All @@ -682,17 +680,17 @@ bool RegBankLegalizeHelper::lowerSBufToBuf(MachineInstr &MI,
default:
break;
}
for (int i = 0; i < NumLoads; ++i) {
for (int I = 0; I < NumLoads; ++I) {
B.buildInstr(Opc)
.addDef(LoadParts[i]) // vdata
.addDef(LoadParts[I]) // vdata
.addUse(RSrc) // rsrc
.addUse(VIndex) // vindex
.addUse(VOffset) // voffset
.addUse(SOffset) // soffset
.addImm(ImmOffset + 16 * i) // offset(imm)
.addImm(0) // cachepolicy, swizzled buffer(imm)
.addImm(ImmOffset + 16 * I) // offset(imm)
.addImm(CachePolicy) // cachepolicy, swizzled buffer(imm)
.addImm(0) // idxen(imm)
.addMemOperand(MF.getMachineMemOperand(BaseMMO, 16 * i, MemSize));
.addMemOperand(MF.getMachineMemOperand(BaseMMO, 16 * I, MemSize));
}
if (NumLoads == 1)
B.buildCopy(Dst, LoadParts[0]);
Expand Down
16 changes: 12 additions & 4 deletions llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -385,7 +385,8 @@ AMDGPURegisterBankInfo::getInstrAlternativeMappingsIntrinsicWSideEffects(
const MachineInstr &MI, const MachineRegisterInfo &MRI) const {

switch (cast<GIntrinsic>(MI).getIntrinsicID()) {
case Intrinsic::amdgcn_s_buffer_load: {
case Intrinsic::amdgcn_s_buffer_load:
case Intrinsic::amdgcn_ptr_s_buffer_load: {
static const OpRegBankEntry<2> Table[4] = {
// Perfectly legal.
{ { AMDGPU::SGPRRegBankID, AMDGPU::SGPRRegBankID }, 1 },
Expand Down Expand Up @@ -1367,9 +1368,9 @@ bool AMDGPURegisterBankInfo::applyMappingSBufferLoad(
LLT Ty = MRI.getType(Dst);

const RegisterBank *RSrcBank =
OpdMapper.getInstrMapping().getOperandMapping(1).BreakDown[0].RegBank;
OpdMapper.getInstrMapping().getOperandMapping(1).BreakDown[0].RegBank;
const RegisterBank *OffsetBank =
OpdMapper.getInstrMapping().getOperandMapping(2).BreakDown[0].RegBank;
OpdMapper.getInstrMapping().getOperandMapping(2).BreakDown[0].RegBank;
if (RSrcBank == &AMDGPU::SGPRRegBank &&
OffsetBank == &AMDGPU::SGPRRegBank)
return true; // Legal mapping
Expand Down Expand Up @@ -1415,6 +1416,7 @@ bool AMDGPURegisterBankInfo::applyMappingSBufferLoad(
Register RSrc = MI.getOperand(1).getReg();
Register VIndex = B.buildConstant(S32, 0).getReg(0);
B.getMRI()->setRegBank(VIndex, AMDGPU::VGPRRegBank);
unsigned CachePolicy = MI.getOperand(3).getImm();

SmallVector<Register, 4> LoadParts(NumLoads);

Expand All @@ -1439,7 +1441,7 @@ bool AMDGPURegisterBankInfo::applyMappingSBufferLoad(
.addUse(VOffset) // voffset
.addUse(SOffset) // soffset
.addImm(ImmOffset + 16 * i) // offset(imm)
.addImm(0) // cachepolicy, swizzled buffer(imm)
.addImm(CachePolicy) // cachepolicy, swizzled buffer(imm)
.addImm(0) // idxen(imm)
.addMemOperand(BaseMMO);
}
Expand Down Expand Up @@ -5464,6 +5466,12 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
OpdsMapping[1] = AMDGPU::getValueMapping(AMDGPU::VCCRegBankID, 1);
break;
}
case Intrinsic::amdgcn_ptr_s_buffer_load: {
OpdsMapping[0] = getSGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI);
OpdsMapping[2] = getSGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
OpdsMapping[3] = getSGPROpMapping(MI.getOperand(3).getReg(), MRI, *TRI);
break;
}
case Intrinsic::amdgcn_raw_buffer_load:
case Intrinsic::amdgcn_raw_ptr_buffer_load:
case Intrinsic::amdgcn_raw_atomic_buffer_load:
Expand Down
Loading
Loading