LG. My preference would be as commented above, but the vector containing the contents of the other two works too.

Thanks!

Yes, I see that you agree with the wording you committed. We can't fix it now.

In D103261#2802471, @hsmhsm wrote:

In D103261#2802464, @JonChesterfield wrote:

I note that you did not change the commit message to reflect the contents of the patch as requested.

The commit message says, what the the patch is doing - it increases the alignment of LDS if it is under-aligned in order to reduce the probability of unlinged access. So, I do not think, any change to commit message is required.

I note that you did not change the commit message to reflect the contents of the patch as requested.

Nice, thanks. The task support was nominally deleted some time ago, iirc the init() part survived because the variable was used elsewhere as well. Good to see it go.

Thanks!

That worked out cleanly, thanks. The temporary vector used to indicate failure with .empty is slightly unusual but obviously works.

Took a while to make sense of this. I think you've preserved the existing strategy of building a collection of kernarg memory pools, while moving the global into a class member and adding some more tests on whether a pool should be added to the collection.

In D103431#2791057, @hsmhsm wrote:

If the current logic of shouldLowerLDSToStruct() is more robust compere to earlier ones, then, why cannot we use it, instead of patching-up earlier ones?

Looks functionally correct. Always in favour of moving stuff out of headers where we can, though as an unrelated change we might want to start building plugins with flto or similar to ensure that remains free.

git-clang-format HEAD^ may be useful for reformatting the already changed parts. LGTM, @jhenderson?

Please continue the review in the existing place (i.e. reopen it) instead of making a new one

If we're going with this anyway, let's be explicit in the title and commit message that we are over-aligning data. We aren't 'fixing' things, or doing them 'properly', we are intentionally burning memory in the hope of improving runtime.

LG, thanks!

Does anyone have a reference for what we're supposed to do with one of these things?

For that test, I note that commenting out the llvm.used statement results in the variable being lowered. I.e. the following works fine:

Is this the proposed fix for the missing lowering of the LDS in the following?

Iirc it did more things in ATMI. Nice simplification here.

I'm not very convinced by this given the recent enthusiasm for decreasing LDS usage. Types are naturally aligned by default, so I think the only time they are aligned by less than that is when the programmer asked for it (or when the vectorizer is involved, but that's disabled on amdgpu afaik).

LGTM, thanks. Cleans up the call sites without changing the calls into HSA. Thanks for dropping the extern C shim too.

Not really the right place for this comment, but better odds of it being seen than in other bug trackers. Mahesha reports:

What's the use case for this? It will increase memory use if it increases the alignment of any variables

We could lower to the equivalent of

void * kernarg_segment_pointer()
{
  __attribute__((address_space(4))) void* p = __builtin_amdgcn_dispatch_ptr();
  void * res;
  __builtin_memcpy(&res, (char*)p + (320/8), 8);
  return res;
}

instead of zero when not in a kernel function. Which is what I've just replaced a call to the builtin with, after a fairly lengthy process of trying to guess why the GPU was reporting memory corruption.

Nice, like that a lot. I hadn't noticed that we are passing device id == 0 around to mean a CPU device, that seems a path to considerable confusion. Probably better to split the allocators into gpu (which takes a device id) and cpu (which doesn't) in some later patch.

Didn't realize we were using string stream. Probably don't want to, but thanks for the fix in the meantime

Ouch. Good to catch that quickly, thanks

• twist code around clang-tidy

• rebase

Use different passes for different things, and optimisations are different from lowering. Those should not be contentious statements in a multipass compiler. Strictly I'm not sure this is an optimisation as it burns cycles to reduce memory, and that's only useful if the saved memory improves occupancy.

I see that. However, we are working on a multipass compiler with a clearly delimited format used between passes.

Same objection to previous diffs with this code in. You're adding loads of complicated code to an existing pass.

LG, thanks!

Thanks for pointing it out! I was under the impression that the contents of 'Summary' as shown in Phabricator turns into the commit message.

passing by const& is much better than passing by &, I think we should do that in this patch

Deleted the references to ATMI_STATUS_UNKNOWN manually, then

for f in `find . -type f`; do sed -i 's$ATMI_STATUS_SUCCESS$HSA_STATUS_SUCCESS$g' $f; sed -i 's$ATMI_STATUS_ERROR$HSA_STATUS_ERROR$g' $f; sed -i 's$atmi_status_t$hsa_status_t$g' $f; done

and clang-format

except for the lost error check above, I think this is now a non-functional change. LG with that error path restored

Was reviewed on aomp, will wait a little while before landing this to match.

Side point on races - I think it is possible for one thread to start loading a binary, get immediately suspended, and a second thread to then launch a kernel that is supposed to come from that binary. The locking scheme in this plugin probably doesn't handle that. Hopefully libomptarget does have locking that makes that work, otherwise the cuda plugin is probably also racy.

In D103030#2778888, @pdhaliwal wrote:

Looks good. My opinion is that we should keep these wrappers in a separate file, but that could be a later patch.

Great, thanks!

Nice, thanks!

Refactoring separately is good, but only works in practice if it's done first. Otherwise the refactor usually gets forgotten, or replaced with something higher priority, and that is the path to entropy.

no objection to the functional change, would be good to refactor the local weirdness in passing

Hostcall exists as a libraries in rocm, aomp and my github. Trunk has hooks - mostly weak symbols - to allow the host plugin source to work with or without it present.

Could change these to std::min

• drop dead executable argument

Couple of minor comments above, overall approach looks sound to me

Like it, thanks!

Couple of style comments above. I think this works as-is, so we have confirmation that the map of strings can go. Which is great!

Ah, apologies. I misread/misremembered this. The global variable looked up here is llvm.used / llvm.compiler.used, and not the LDS variable that is going to be removed.

That is much prettier, thank you.

I'm expecting to have a few in flight locally while trying different designs. I don't expect this to be a long lasting patch.

In D94648#2770088, @hsmhsm wrote:

Hi Jon,

For the 0 argument kernel, did you put debug statement within SITargetLowering::LowerFormalArguments() and test whether it will hit or not? My experimentation shows that it is indeed hit for 0 arg kernels too. So it is not problem with 0 arg kernel.

Taken a closer look at this. Cut down example msgpack,

".name" : "__omp_vmul_l8",
".symbol" : "__omp_vmul_l8.kd",