We read every piece of feedback, and take your input very seriously.
To see all available qualifiers, see our documentation.
There was an error while loading. Please reload this page.
1 parent b14f1e5 commit c4b2b8bCopy full SHA for c4b2b8b
llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
@@ -45,6 +45,14 @@ bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) {
45
return true;
46
}
47
48
+// In HIP, dynamic LDS variables are represented using 0-element global arrays
49
+// in the __shared__ language address-space.
50
+//
51
+// extern __shared__ int LDS[];
52
53
+// These are not representable in SPIRV directly.
54
+// To represent them, for AMD, we use an array with UINT32_MAX-elements.
55
+// These are reverse translated to 0-element arrays.
56
bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) {
57
constexpr unsigned WorkgroupAS =
58
storageClassToAddressSpace(SPIRV::StorageClass::Workgroup);
0 commit comments