yeah, here one example

inttoptr and ptrtoint should be treated as opaque ones. If needed, they should be handled with target-specific. This pass is definitely not the place to handle them.

revise after review comments

revise after review comments.

In D72709#1820182, @rampitec wrote:

You have skipped the dead MO, but was pass reordering really necessary? It seems we have higher register pressure with this change.

In D72709#1820225, @arsenm wrote:

I'm not sure this logic works in the presence of an undef subregister def

In D72174#1803485, @arsenm wrote:

Where is this being called from that it's so expensive?

one compilation is reduced from 460s to 121s after this trivial change.

Could you double-check that on other platforms? Like Linux with shared library builds. At least, there are crash in my local build.

code refinement after reviewers' comments.

refinements are made after comments from reviewers.

Add clarification on the test case.

Add test case.

In D71227#1778136, @tra wrote:

I wonder if this patch will help with this case:

https://godbolt.org/z/X4KdsV

__device__ float fn(int) { return threadIdx.x; };
__host__ float fn(float);

float gvar1 = []()__device__ { return fn(1);} (); // This ends up calling fn(int) on *host*

We seem to happily let host code call device function from a lambda function used as an initializer.

File PR44266 to track that bug.

refine again

refine commit message

LGTM, anyway calling a virtual function inside the constructor is not encouraged.

revise a test

This patch is revived with more changes addressing the previous concerns.

PING for review

PING for review

Sam, could you review this? Even though it has no functionality issue so far, from the code sequence, once there's an addrspacecast is inserted, we lose the chance to set target specific attributes if any.

In D69826#1734296, @yaxunl wrote:

I am a little bit concerned that user may have such code:

struct A { int *p; }
__global__ kernel(A a) {
  int x;
  a.p = &x;
  f(a);
}

@arsenm what happens if a private pointer is mis-used as a global pointer?

I am wondering if we should coerce byval struct kernel arg to global only if they are const, e.g.

__global__ kernel(const A a);

I understand this may lose performance. Or should we introduce an option to let user disable coerce of non-const struct kernel arg to global.

A different approach (D69826) is taken to address the same issue.

Add host-side checks.

revise parameter names

• revise member function name.
• add the test case for by-val array types.

revise code following reviwers' comments.

add the test case for struct.

It happens that Sam has a similar patch of this one. After discussion, we agreed that this patch addresses more cases found in the workloads. Thank Sam for the test case.

Looks like holidays are approaching, :). Anyway, it's really appreciated that you could review this patch to enable CUDA/HIP applications on Windows.

In D69735#1730811, @arsenm wrote:

What was the symptom of this?

could you keep the original scoped lock by adding a level of block?

kindly PING for review

simplify again following suggestion.

revise MSHIPNumberingContext

PING for review