Fix typo.

Forget that C function could be overloaded on Clang with overloadable
extension. With that, we don't need to mark functions from <ymath.h> as HD.
Instead, we could provide their device-side implementation directly.

Only mark HD attributes in ymath.h wrapper header when compiled with MSVC.

Revise following reviewers' comments.

PING

PING

Fix the cmake to distribute that header wrapper.

These functions are pure C functions.

In D93638#2466097, @jdoerfert wrote:

Disclaimer: I request changes because of the next sentence, other than that I have no objection but also cannot review this.
All cuda_wrapper headers say something about complex in the first row, copy & paste error. All have the wrong license text (I think).

Fix license.

Fix typo again.

Fix typo.

Beyond the enabling of the compilation with <complex> on Windows, I really have the concern on the current approach supporting <complex> compilation in the device compilation. The device compilation should not relies on the host STL implementation. That results in inconsistent compilation results across various platforms, especially Linux vs. Windows.
BTW, the use of <complex> in CUDA cannot be compiled with NVCC directly even with --expt-relaxed-constexpr, c.f. https://godbolt.org/z/3f79co

The build is broken due to the missing file.

LGTM if you revise the test based on Sam's suggestion on the test case.

LGTM if there's a regression test available.

Even there's no functionality change, the original one breaks the kernel extraction script, which is designed to find the .hip_fatbin section. That internal tool is still required to extract kernels from objects generated from RDC linking.

In D91928#2410283, @jlebar wrote:

In case it's not used in PHI or SELECT and cannot ensure the result is also a pointer to the parameter space, we could skip alloca insertion.

I think an allowlist might be more appropriate than a denylist. Rather than, anything other than PHI and SELECT, could it be, if it's only transitively used by gep and load we're good?

I am not 100% sure even that works, though. The real problem is that this pass is trying to reason about what the addrspace inference pass is capable of. We can only do the transformation if here if we're positive that addrspace inference will eliminate all generic loads from the arg. That's a layering violation and ultimately is fragile.

In D91928#2410164, @jlebar wrote:

I don't believe there's any exception to prove deduction [of the readonly attribute] wrong.

Understood.

The address space inference here only refers to the one in the backend directly after this argument lowering gpass.

Also understood.

This isn't speaking to my concern, though.

Suppose we have

__global__ void foo(int x, const int* y, int* out, bool flag) {
  int* ptr = flag ? &x : y;
  *out = *ptr;
}

In this case we can say with confidence that x is readonly.

But address space inference cannot infer the address space of ptr (how could it?). Therefore we will do a generic load, which is wrong.

In D91928#2409971, @jlebar wrote:

This looks really simple, which is awesome. I am enthusiastic. But I am worried it may not be correct.

AIUI params are special in that they *must* be read from the param address space. It is illegal to do a generic load of a param.

So this change is correct only if we can guarantee that address space inference will infer the specific address space for all uses of the pointer.

But address space inference is not guaranteed. For example, you could select on two pointers of two different address spaces. So long as you only ever read from these pointers, the arg can still be marked as ReadOnly. But with this patch, we'd end up doing a generic load from the param space, which would be illegal.

Take it all with a grain of salt since I've also been out of the game for a while.

It turns out that the simplest way is to skip generating alloca once that byval argument is readonly. As readonly will be attributed once there's no write to that argument, it's safe to just cast that pointer to the parameter space if it has readonly. Basically, that argument lowering pass does a similar to D91590 but, instead, applies that in the backend. I verified that, for that simple test CUDA code, it would generate the same SASS.

Do you have permission to commit?

In D91513#2402005, @guopeilin wrote:

In D91513#2401715, @hliao wrote:

In D91513#2401473, @guopeilin wrote:

In D91513#2401470, @guopeilin wrote:

In D91513#2400468, @hliao wrote:

Using post-order is quite straight-forward and only involves several lines of change. Please check the attachment.

0001-DeadMachineInstructionElim-Use-post-order-for-basic-.patch4 KBDownload

That test passed with this traverse order change.

That's a great help, I pass all my related cases with this patch, Thanks a lot.

Now that we decide to use post order to visit all blocks of a function, I think we need to consider that what if CFG contains cycles?

From this picture, we can see that post order is not clearly defined cause there exits cycles, one of the possible orders is that [ m, g, d, e, c, b, t, x]
So m comes before g, if we define something in m and use it in g. Then even though both def and use are useless, cause we visit m first, we will still get a dead definition after we post-order visit all blocks.
So is it possible there still exist some cases theoretically that cannot be fixed by post-order visit? That is we may still need to iteratively run?

You are right, that's possible. That case should be rare as that's a def in the back-edge with acyclic dep. Could you merge the post-order change together with the iterative runs? so that, in the regular case, we at most run twice. Please keep on eye on compile time.

Post order visit and iteratively run are merged.

As mentioned earlier, that's very experimental support. Even though the SASS looks reasonable, it still needs verifying on real systems. For non-kernel functions, it seems we share the path. So that we should do a similar thing. The current approach fixes that in the codegen phase by adding back the alloca to match the parameter space semantic. Once that alloca is dynamically indexed, it won't be promoted in SROA. Only instcomb eliminates that alloca when it is only modified once by copying from a constant memory. As instcomb won't break certain patterns prepared in the codegen preparation, it won't run in the backend. That dynamically indexed alloca won't be removed.

BTW, please add a test case with that def in back-edge with acyclic dep.

In D91513#2401473, @guopeilin wrote:

In D91513#2401470, @guopeilin wrote:

In D91513#2400468, @hliao wrote:

Using post-order is quite straight-forward and only involves several lines of change. Please check the attachment.

0001-DeadMachineInstructionElim-Use-post-order-for-basic-.patch4 KBDownload

That test passed with this traverse order change.

That's a great help, I pass all my related cases with this patch, Thanks a lot.

Now that we decide to use post order to visit all blocks of a function, I think we need to consider that what if CFG contains cycles?

From this picture, we can see that post order is not clearly defined cause there exits cycles, one of the possible orders is that [ m, g, d, e, c, b, t, x]
So m comes before g, if we define something in m and use it in g. Then even though both def and use are useless, cause we visit m first, we will still get a dead definition after we post-order visit all blocks.
So is it possible there still exist some cases theoretically that cannot be fixed by post-order visit? That is we may still need to iteratively run?

Using post-order is quite straight-forward and only involves several lines of change. Please check the attachment.

This's an experimental or demo-only patch in my spare time on eliminating private memory usage in https://godbolt.org/z/EPPn6h. The attachment

In D91513#2398733, @rampitec wrote:

In D91513#2398726, @hliao wrote:

In D91513#2398717, @rampitec wrote:

In D91513#2398697, @hliao wrote:

In D91513#2398528, @guopeilin wrote:

In D91513#2397313, @hliao wrote:

could you elaborate more on why we need to run that iteratively? since the original one runs bottom-up, supposedly it should find all.

From the iteratively-run-dead-mi-elim.mir we can see that bb.5 defines %6, and %6 is used in bb.2. When we traverse the all basic blocks, that is, we runs bottome-up, we will meet bb.5 first, for %6, we find that it is not dead cause %3 in bb.2 use it. So %6 surrive. Then we continue traverse other BBs, When we meet bb.2, we see that no one use %5, so we kill it. So as %4, %3. Right now, actually %6 becomes dead cause we kill %3 thus there is no longer any one uses %6.
However, cause we only traverse blocks once, we can't erase %6 at the end. So if we iteratively visit all blocks until nothing change, then we can ensure that all dead mi is erased.

Ah, I see. Even though we try to traverse basic blocks bottom-up, that's just the block placement order instead of block reachability. Could we replace that order with the post-order? So that, the use is always traversed before the defiine.

That probably will not help if we have a loop?

It still works unless that value has a cyclic dependency through phi-node.

That's exactly what I had in mind, a phi node as the only way to get a cyclic dependency in SSA.

I tend to say this is LGTM. Although I wish to see a test with a cyclic dependency.

In D91513#2398717, @rampitec wrote:

In D91513#2398697, @hliao wrote:

In D91513#2398528, @guopeilin wrote:

In D91513#2397313, @hliao wrote:

could you elaborate more on why we need to run that iteratively? since the original one runs bottom-up, supposedly it should find all.

From the iteratively-run-dead-mi-elim.mir we can see that bb.5 defines %6, and %6 is used in bb.2. When we traverse the all basic blocks, that is, we runs bottome-up, we will meet bb.5 first, for %6, we find that it is not dead cause %3 in bb.2 use it. So %6 surrive. Then we continue traverse other BBs, When we meet bb.2, we see that no one use %5, so we kill it. So as %4, %3. Right now, actually %6 becomes dead cause we kill %3 thus there is no longer any one uses %6.
However, cause we only traverse blocks once, we can't erase %6 at the end. So if we iteratively visit all blocks until nothing change, then we can ensure that all dead mi is erased.

Ah, I see. Even though we try to traverse basic blocks bottom-up, that's just the block placement order instead of block reachability. Could we replace that order with the post-order? So that, the use is always traversed before the defiine.

That probably will not help if we have a loop?

In D91513#2398528, @guopeilin wrote:

In D91513#2397313, @hliao wrote:

could you elaborate more on why we need to run that iteratively? since the original one runs bottom-up, supposedly it should find all.

From the iteratively-run-dead-mi-elim.mir we can see that bb.5 defines %6, and %6 is used in bb.2. When we traverse the all basic blocks, that is, we runs bottome-up, we will meet bb.5 first, for %6, we find that it is not dead cause %3 in bb.2 use it. So %6 surrive. Then we continue traverse other BBs, When we meet bb.2, we see that no one use %5, so we kill it. So as %4, %3. Right now, actually %6 becomes dead cause we kill %3 thus there is no longer any one uses %6.
However, cause we only traverse blocks once, we can't erase %6 at the end. So if we iteratively visit all blocks until nothing change, then we can ensure that all dead mi is erased.

Kindly ping for review.

could you elaborate more on why we need to run that iteratively? since the original one runs bottom-up, supposedly it should find all.

Revise the interface of that target hook.
Add a dedicated test case for value reading from parameter even though most cases are already covered in the clang test.

Revise the condition check.

• Add a note in the AMDGPU usage document on the assumption made here.
• Revise the test in clang.

Add a test case for the single element struct.

PING for review

PING for review

Rebase

Fix clang-tidy warnings.

Revise the fix.

Revise the commit message.

Remove aggregate kernel argument coercion only.

with multiple MMO is supported in the scheduler, this patch is no longer for performance.