- HIP uses an unsized extern array extern __shared__ T s to declare the dynamic shared memory, which size is not known at the compile time.
Needs to handle globalisel too.
I also thought we were trying to get rid of group static size. It's broken with LDS relocations which we need to move towards. Can we just switch directly to using a relocation here?
Why does it specifically need to be a 0 sized array? I think this would depend purely on the linkage, or treat any 0 sized type the same way
I just found that change for non-HSA/-PAL environment. I need to check how it works and fit into other tests. So far, that's a critical change to ensure we won't change the original source code too much. Is it possible to address that relocation in a long run (says 1~3 weeks) to avoid the tight schedule.
That's the only syntax accepted under HIP clang. At least, zero-sized types should be checked to ensure developers understand the usage of the dynamic shared memory.
My understanding is this feature is equivalent to the OpenCL dynamic group segment allocation. The runtime would presumably implement it in a similar way.
So the HIP runtime must take the static LDS size, round up to the alignment requirement of the dynamic allocation (OpenCL just uses the maximally aligned OpenCL data type), then add the size of the dynamic LDS. The AQL packet group segment field is set to the total LDS size.
In OpenCL there can be multiple kernel arguments, and the LDS address is passed to each. But for HIP there is only one dynamic area denoted by this weird extern. How is the dynamic LDS storage accessed? Is the address passed as an implicit kernel argument, or does the compiler implicitly use the aligned static LDS size?
I don't think this actually works since you could have multiple 0 sized objects, and they would both get the same address. I think this has to be an external relocation
That's a bit too much HIP specific logic. Also what does this do if there are more than one? How can these return different addresses?
The global has an explicit alignment that needs to be respected
This's the point. To keep compatible with CUDA, multiple dynamically sized arrays in a single kernel must declare a single extern unsized array and uses the address to divide it into multiple arrays by developers themselves. It's in fact quite similar to the local memory parameter in OpenCL. Thus, all extern unsized __shared__ arrays are mapped onto the same address. The developer has the responsibility to divide it into multiple ones. That's in fact is same across HCC and HIP-Clang except that we want to maximize the compatibility and avoid changing source code too much.
Also missing the globalisel handling
I don't think this should be mutable
This doesn't need to be MaybeAlign, just Align. Also expand to Dynamic?
Also should probably elaborate that this is used for the case where a dynamically sized global is used
Rename to getAllocatedLDSSize()? I think there should be a separate method to get the size plus the roundup to the dynamic alignment
Set is the wrong word here; ensureDynamicLDSAlign()?
This should check if the allocated size of GV->getValueType() is 0, not special case 0 sized arrays
Should be the alignment of the aggregate itself, not the element type
To report the LDS usage accurately, the final LDSSize need to count the padding due to dynamic shared memory alignment. Now, the re-alignment on LDS is explicitly done before the final instruction selection.
The re-alignment is done explicitly just before final instruction selection and just once. BTW, getLDSSize should be only valid after instruction selection.
I think these should remain distinct queries/fields, not fixed up at an arbitrary point. GlobalISel will miss this for example. The asm printer would query the kind that accounts for the dynamic padding
This should not special case 0 sized arrays and should check the allocated size
This can still just be Align
This isn't a set since it does more than set the value. ensureDynamicLDSAlign?
This shouldn't be a mutation, but return the aligned up size.
This should not special case 0 sized arrays. This is 0 allocation size
The default should be 1
GlobalISel calls adjustLDSSizeForDynLDSAlign similarly before the finalization of ISel.
I tend to be restrictive here to follow how that is used in HIP, zero-sized array always has zero allocated size. If other clients need similar usage but general zero-allocated type, we may enhance accordingly.
don't we need to report the accurate usage of LDS? Does that alignment padding need counting as well for the final LDSSize?
Having extra state that needs to be made consistent is bad. It's better to just track the two independent fields and do the roundup when needed at the end
Types don't mean anything. Any 0 sized globals are getting allocated to the same address. We're just going to miscompile other 0 sized types. We have no reason to treat other 0 sized types differently
Should use the aggregate alignment, not the element