Page MenuHomePhabricator

[HIP] change kernel stub name
ClosedPublic

Authored by yaxunl on Feb 21 2019, 10:02 AM.

Diff Detail

Repository
rC Clang

Event Timeline

yaxunl created this revision.Feb 21 2019, 10:02 AM
tra accepted this revision.Feb 21 2019, 11:07 AM

My guess is that this is needed because HIP debugger can see symbols from both host and device executables at the same time. Is that so?

If that's the case, I guess HIP may have similar naming problem for __host__ __device__ foo() if it's used on both host and device.

lib/CodeGen/CGCUDANV.cpp
230–231

It may be worth adding a comment why kernel stub in HIP needs a different name.

This revision is now accepted and ready to land.Feb 21 2019, 11:07 AM
yaxunl marked an inline comment as done.Feb 21 2019, 11:24 AM
In D58518#1406124, @tra wrote:

My guess is that this is needed because HIP debugger can see symbols from both host and device executables at the same time. Is that so?

If that's the case, I guess HIP may have similar naming problem for __host__ __device__ foo() if it's used on both host and device.

Probably, will fix it in seperate patch if it is true.

lib/CodeGen/CGCUDANV.cpp
230–231

will do when commit

t-tye added a comment.EditedFeb 21 2019, 11:37 AM

Yes this relates to supporting the debugger.

For the same function being present on both host and device, having the same name is correct as the debugger must set a breakpoint at both places. This is similar to needing to set a breakpoint at every place a function is inlined.

tra added a comment.Feb 21 2019, 11:55 AM

Yes this relates to supporting the debugger.

For the same function being present on both host and device, having the same name is correct as the debugger must set a breakpoint at both places. This is similar to needing to set a breakpoint at every place a function is inlined.

I'm confused. Are you saying that HIP does *not* need a different name for the stub then?

To clarify, I am saying that the stub does have a different name since it is conceptually part of the implementation of doing the call to the device function implementation, and is not in fact the the device function being called itself. However, when we generate code for a function that is present on both the host and device, both copies of the code are for the same source level function and so can have the same symbol name (which was a question that was asked).

Closed by commit rC354615: [HIP] change kernel stub name (authored by yaxunl, committed by ). · Explain WhyFeb 21 2019, 12:11 PM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptFeb 21 2019, 12:11 PM
tra added a comment.Feb 21 2019, 1:01 PM

To clarify, I am saying that the stub does have a different name since it is conceptually part of the implementation of doing the call to the device function implementation, and is not in fact the the device function being called itself. However, when we generate code for a function that is present on both the host and device, both copies of the code are for the same source level function and so can have the same symbol name (which was a question that was asked)

Got it. Agreed.

yaxunl updated this revision to Diff 187980.Feb 22 2019, 1:44 PM

Fixed regressions.

yaxunl reopened this revision.Feb 22 2019, 1:45 PM
This revision is now accepted and ready to land.Feb 22 2019, 1:45 PM
tra requested changes to this revision.Feb 22 2019, 2:20 PM
tra added a subscriber: echristo.
tra added inline comments.
lib/CodeGen/CodeGenModule.cpp
1059 ↗(On Diff #187980)

Changing mangled name exposes this change to a wider scope of potential issues. Is the mangled name still valid after this change? I.e. will external demanglers have problem with it? Is . a valid symbol in mangled names on all platforms we support?

I think changing the name here is way too late and we should figure out how to change the stub name when we generate it.

@echristo Eric, what do you think?

This revision now requires changes to proceed.Feb 22 2019, 2:20 PM
yaxunl added inline comments.Feb 22 2019, 3:46 PM
lib/CodeGen/CodeGenModule.cpp
1059 ↗(On Diff #187980)

The external demangler can still demangle this name. e.g. c++filt will demangle this name and add [clone .stub] after that.

As far as I can see this function is only called in codegen to map FunctionDecl names to LLVM function names. I've tried this change with real ML frameworks and it works.

Changing at this place is not too late. The stub function name is requested at multiple places in codegen, not just at the emitting of stub function definition. For template kernel function, the emitting of stub function definition is deferred after emitting of the call of the stub function. Basically, codegen needs to find the corresponding LLVM stub function by getMangledName first, then by GetOrCreateLLVMFunction. If we do not change getMangledName, codegen will not get the correct stub function name consistently at all places. That's why the previous patch does not work.

tra accepted this revision.Feb 26 2019, 2:05 PM
tra added subscribers: jyknight, bkramer.
tra added inline comments.
lib/CodeGen/CodeGenModule.cpp
1059 ↗(On Diff #187980)

I stand corrected. @jyknight and @bkramer pointed out that appending .WHATEVER is currently used for cloning functions and should be OK to do.

This revision is now accepted and ready to land.Feb 26 2019, 2:05 PM
Closed by commit rL354948: [HIP] change kernel stub name (authored by yaxunl, committed by ). · Explain WhyFeb 26 2019, 6:03 PM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptFeb 26 2019, 6:03 PM