Skip to content

[SPIRV][SPIRVPrepareGlobals] Map AMD's dynamic LDS 0-element globals to arrays with UINT32_MAX elements #166952

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

Merged
merged 5 commits into from
Nov 12, 2025
Merged

[SPIRV][SPIRVPrepareGlobals] Map AMD's dynamic LDS 0-element globals to arrays with UINT32_MAX elements #166952

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
89f8cc5
[SPIRV][SPIRVPrepareGlobals] Map AMD's dynamic LDS 0-element globals …
jmmartinez Nov 7, 2025
4bccb7f
[Review] Rename Magic->UInt32Max
jmmartinez Nov 10, 2025
b14f1e5
[Review] use storageClassToAddressSpace(SPIRV::StorageClass::Workgroup);
jmmartinez Nov 10, 2025
c4b2b8b
[Review] Add comment
jmmartinez Nov 10, 2025
ab31a73
Merge branch 'main' into users/jmmartinez/spirv/hip_dyn_lds
jmmartinez Nov 12, 2025
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
37 changes: 37 additions & 0 deletions llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,9 @@
//===----------------------------------------------------------------------===//

#include "SPIRV.h"
#include "SPIRVUtils.h"

#include "llvm/ADT/STLExtras.h"
#include "llvm/IR/Module.h"

using namespace llvm;
Expand Down Expand Up @@ -43,6 +45,38 @@ bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) {
return true;
}

// In HIP, dynamic LDS variables are represented using 0-element global arrays
// in the __shared__ language address-space.
//
// extern __shared__ int LDS[];
//
// These are not representable in SPIRV directly.
// To represent them, for AMD, we use an array with UINT32_MAX-elements.
// These are reverse translated to 0-element arrays.
bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) {
constexpr unsigned WorkgroupAS =
storageClassToAddressSpace(SPIRV::StorageClass::Workgroup);
const bool IsWorkgroupExternal =
GV.hasExternalLinkage() && GV.getAddressSpace() == WorkgroupAS;
if (!IsWorkgroupExternal)
return false;

const ArrayType *AT = dyn_cast<ArrayType>(GV.getValueType());
if (!AT || AT->getNumElements() != 0)
return false;
Comment on lines +64 to +66
Copy link
Contributor

@s-perron s-perron Nov 10, 2025

Choose a reason for hiding this comment

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

What do you want to do with 0-sized arrays that are not the type of the global value? Is even possible to do that? Comments explaining why you limit this to just the type of the GV would be useful.

Copy link
Contributor Author

@jmmartinez jmmartinez Nov 10, 2025

Choose a reason for hiding this comment

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

Sorry I'm not sure I understood the question.

This condition is matching globals that have an array type with 0 elements. If the global does not have an array type or if it is an array type with a size different from 0 this function returns false.

Maybe there is a mix from using getValueType vs getType. The first gives the type of the initializer of the global (the value stored in the global), the second returns the type of the global when used as a value in the llvm-ir (a pointer type).

Copy link
Contributor

@s-perron s-perron Nov 10, 2025

Choose a reason for hiding this comment

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

Sorry, I did not write that properly.

Could you have, say, a global whose type is a struct containing a 0-sized array? What do you want to do in case?

@lds = external addrspace(3) global {i32, [0 x i32]}

Copy link
Contributor Author

@jmmartinez jmmartinez Nov 12, 2025

Choose a reason for hiding this comment

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

I don't think it's possible. These cases get rejected by the frontend normally: https://godbolt.org/z/8PdqqrYMT


constexpr auto UInt32Max = std::numeric_limits<uint32_t>::max();
ArrayType *NewAT = ArrayType::get(AT->getElementType(), UInt32Max);
GlobalVariable *NewGV = new GlobalVariable(
*GV.getParent(), NewAT, GV.isConstant(), GV.getLinkage(), nullptr, "",
&GV, GV.getThreadLocalMode(), WorkgroupAS, GV.isExternallyInitialized());
NewGV->takeName(&GV);
GV.replaceAllUsesWith(NewGV);
GV.eraseFromParent();

return true;
}

bool SPIRVPrepareGlobals::runOnModule(Module &M) {
const bool IsAMD = M.getTargetTriple().getVendor() == Triple::AMD;
if (!IsAMD)
Expand All @@ -52,6 +86,9 @@ bool SPIRVPrepareGlobals::runOnModule(Module &M) {
if (GlobalVariable *Bitcode = M.getNamedGlobal("llvm.embedded.module"))
Changed |= tryExtendLLVMBitcodeMarker(*Bitcode);

for (GlobalVariable &GV : make_early_inc_range(M.globals()))
Changed |= tryExtendDynamicLDSGlobal(GV);

return Changed;
}
char SPIRVPrepareGlobals::ID = 0;
Expand Down
20 changes: 20 additions & 0 deletions llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
; RUN: llc -verify-machineinstrs -mtriple=spirv64-amd-amdhsa %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -mtriple=spirv64-amd-amdhsa %s -o - -filetype=obj | spirv-val %}

; CHECK: OpName %[[#LDS:]] "lds"
; CHECK: OpDecorate %[[#LDS]] LinkageAttributes "lds" Import
; CHECK: %[[#UINT:]] = OpTypeInt 32 0
; CHECK: %[[#UINT_MAX:]] = OpConstant %[[#UINT]] 4294967295
; CHECK: %[[#LDS_ARR_TY:]] = OpTypeArray %[[#UINT]] %[[#UINT_MAX]]
; CHECK: %[[#LDS_ARR_PTR_WG:]] = OpTypePointer Workgroup %[[#LDS_ARR_TY]]
; CHECK: %[[#LDS]] = OpVariable %[[#LDS_ARR_PTR_WG]] Workgroup

@lds = external addrspace(3) global [0 x i32]

define spir_kernel void @foo(ptr addrspace(4) %in, ptr addrspace(4) %out) {
entry:
%val = load i32, ptr addrspace(4) %in
%add = add i32 %val, 1
store i32 %add, ptr addrspace(4) %out
ret void
}
Loading