Page MenuHomePhabricator

[HIP] Support attribute hip_pinned_shadow
ClosedPublic

Authored by yaxunl on May 31 2019, 8:35 AM.

Details

Summary

This patch introduces support of hip_pinned_shadow variable for HIP.

A hip_pinned_shadow variable is a global variable with attribute hip_pinned_shadow.
It has external linkage on device side and has no initializer. It has internal
linkage on host side and has initializer or static constructor. It can be accessed
in both device code and host code.

This allows HIP runtime to implement support of HIP texture reference.

Diff Detail

Repository
rL LLVM

Event Timeline

yaxunl created this revision.May 31 2019, 8:35 AM
tra added a comment.Jun 11 2019, 11:30 AM

Syntactically the patch looks OK to me, but I think the purpose and meaning of the builtin type should be documented in more details. Based on this patch alone it's not clear at all what it's supposed to be used for and how.

include/clang/Basic/AttrDocs.td
4166 ↗(On Diff #202440)

What does it mean for user-defined template to be a builtin type? This sounds like contradiction to me.

At the very least it should be documented somewhere what are the requirements for such a template and what is it supposed to do in the end.

lib/CodeGen/CodeGenModule.cpp
2424 ↗(On Diff #202440)

Nit: I'd check LangOpts.HIP first. No need chasing pointers in isCUDATExtureType if it's not a HIP compilation.

3781 ↗(On Diff #202440)

This is not predicated by HIP compilation and will have effect on CUDA.

It's also not clear to me why texture is initialized as undef on device side. Adding a comment about that would be great.

yaxunl updated this revision to Diff 204155.Jun 11 2019, 1:42 PM
yaxunl marked 5 inline comments as done.

Revised by Artem's comments.

tra added a comment.Jun 11 2019, 2:22 PM

So, the only thing this patch appears to do is make everything with this attribute uninitialized on device side and give protected visibility.
If I understand it correctly, you're using the attribute in order to construct something that's sort of opposite of the currently used device vars with host-side shadows. Only now the real variable lives on the host side and it's device side that gets the 'shadow' copy. Do I understand it correctly?

If so, then this functionality seems to be fairly general-purpose to me. I.e. it has literally nothing to do with textures other than the name.

Perhaps it would make more sense to rename this attribute to something along the lines of 'device_referenceable' and bring its implementation to somewhat more complete shape.

By 'complete' I mean that it would be great to flesh out what can and can't use the attribute. Does it have to be a type attribute, or can it be applied to variables?
The example in the patch suggests that it's the *variable* that's affected by the attribute.

Once it works, HIP's texture support can use it for its purposes.

E.g. your example could look like this:

#define  __attribute__((device_builtin_texture_type)) __texture__

template <class T, int texType, enum hipTextureReadMode>
    struct  texture
      : public textureReference { ... }

__texture__ texture<float, 2, hipReadModeElementType> tex;

This way compiler does not need to deal with the details of texture implementation on the HIP side.
Host/device visibility of the variables is easy to see in the source (similar to device, shared, etc) and there will be no need to dig into template defined somewhere else to become aware of this.
It will be potentially useful beyond HIP-only texture implementation.

What do you think?

tra added a reviewer: jlebar.Jun 11 2019, 2:23 PM
yaxunl updated this revision to Diff 205490.Jun 18 2019, 6:44 PM
yaxunl retitled this revision from [HIP] Support texture type to [HIP] Support device_shadow variable.
yaxunl edited the summary of this revision. (Show Details)

Revised by Artem's comments.

yaxunl updated this revision to Diff 205503.Jun 18 2019, 9:33 PM

Fix visibility and dso_local. Allow undefined symbol in code object. This is to allow merging the host and device symbols at run time.

tra added a comment.Jun 19 2019, 11:05 AM

Overall looks good. Thank you for making the change.

While reviewing the patch it occured to me that it presents an opportunity to generalize the shadow variables to work in both directions. See below.

include/clang/Basic/AttrDocs.td
4164–4171 ↗(On Diff #205503)

just device shadow variable would do. It's no longer, generally speaking, HIP-specific. :-)

Only address and size of such variables should be used on device side.

I'd rephrase the use constraint. Currently it's !(CUDA || !CUDA) which is always false.
Currently enabled for HIP only. would be closer to reality.

lib/CodeGen/CodeGenModule.cpp
3775 ↗(On Diff #205503)

IsDeviceShadowVar. We may want to rename IsCUDAShadowVar to IsHostShadowVar to be consistent.

This got me thinking. Conceptually we have two different things going on here.

  • where do we place the real variable
  • whether we need to create a shadow on the other end.

Currently __device__, __constant__ and __shared__ act as both.
This patch implements the same make a shadow on the other side, only in the opposite direction.

Perhaps the right thing to do is to push the patch even further and make it into a __shadow_variable__ which will be responsible for creating the other side shadow and would work in both directions.

We can then assign implicit __shadow_variable__ attribute to the device-side vars to preserve current behavior and it will work for your purposes two. We will also gain ability to create device-side variables w/o host-side shadows, if we need to.

I guess in the end it would be this patch + a bit of refactoring/collapsing of IsCUDAShadowVar logic.

yaxunl marked an inline comment as done.Jun 20 2019, 9:36 AM
In D62738#1538900, @tra wrote:

So, the only thing this patch appears to do is make everything with this attribute uninitialized on device side and give protected visibility.
If I understand it correctly, you're using the attribute in order to construct something that's sort of opposite of the currently used device vars with host-side shadows. Only now the real variable lives on the host side and it's device side that gets the 'shadow' copy. Do I understand it correctly?

If so, then this functionality seems to be fairly general-purpose to me. I.e. it has literally nothing to do with textures other than the name.

Perhaps it would make more sense to rename this attribute to something along the lines of 'device_referenceable' and bring its implementation to somewhat more complete shape.

By 'complete' I mean that it would be great to flesh out what can and can't use the attribute. Does it have to be a type attribute, or can it be applied to variables?
The example in the patch suggests that it's the *variable* that's affected by the attribute.

Once it works, HIP's texture support can use it for its purposes.

E.g. your example could look like this:

#define  __attribute__((device_builtin_texture_type)) __texture__

template <class T, int texType, enum hipTextureReadMode>
    struct  texture
      : public textureReference { ... }

__texture__ texture<float, 2, hipReadModeElementType> tex;

This way compiler does not need to deal with the details of texture implementation on the HIP side.
Host/device visibility of the variables is easy to see in the source (similar to device, shared, etc) and there will be no need to dig into template defined somewhere else to become aware of this.
It will be potentially useful beyond HIP-only texture implementation.

What do you think?

The problem is that we do not see generic usage of
Although there is no texture specific handling on the compiler side, there is texture specific handling of symbols

include/clang/Basic/AttrDocs.td
4164–4171 ↗(On Diff #205503)

If only address and size of such variables should be used on device side, such variables will not be very useful.

To implement texture reference, we need to be able to load the device side shadow variable. In general, it is desirable to load and store device side shadow variables, since users have no other way to synch with the corresponding host variable in device code.

This is different from host side shadow variable. On host side, users can use hipMemcpyToSymbol and hipMemcpyFromSymbol to force synchronization between the host side shadow variable and the corresponding device variable.

Therefore the implementation of the device side shadow variable requires special handling in HIP runtime. Basically HIP runtime needs to pin the host variable and use it to resolve the device side shadow variable (as an undefined elf symbol). This way, the host variable and device side shadow variable are sharing the same memory. This is also why it is HIP specific since CUDA runtime may not have such handling.

The problem is that we do not see generic usage of
Although there is no texture specific handling on the compiler side, there is texture specific handling of symbols

Please ignore this comment. It is some old comment submitted by accident.

yaxunl marked an inline comment as done.Sun, Jun 23, 2:24 PM
yaxunl added inline comments.
lib/CodeGen/CodeGenModule.cpp
3775 ↗(On Diff #205503)

Do we really want to introduce a generic __shadow_variable__ for device variables? It has little use but complicates AST of device variables unnecessarily. First it bring no new functionality since device variables are already shadowed by default. Second since unused shadow variable is eliminated automatically due to their internal linkage, disable shadowing will not save memory in host binary.

tra added inline comments.Mon, Jun 24, 10:38 AM
include/clang/Basic/Attr.td
954 ↗(On Diff #205503)

HIPDeviceShadow ?

955 ↗(On Diff #205503)

In light of the details you've provided below, perhaps this needs a better name. I've suggested __device_shadow__ without being aware of what exactly it's supposed to do in HIP.
Perhaps something like __hip_device_shadow__ or __hip_pinned_shadow ? Naming is hard. :-)

957 ↗(On Diff #205503)

Shis should probably be [HIP] now, too.

include/clang/Basic/AttrDocs.td
4164–4171 ↗(On Diff #205503)

Thank you for providing the details. This use case is sufficiently different from the more general purpose reverse-shadow-var I had in mind.

lib/CodeGen/CodeGenModule.cpp
3775 ↗(On Diff #205503)

There's currently no specific use case for it in CUDA and HIP's use case also does not quite fit this, so __shadow_variable__ is not going to do either of us any good. So, we can drop the __shadow_variable__.

yaxunl marked an inline comment as done.Mon, Jun 24, 11:56 AM
yaxunl added inline comments.
include/clang/Basic/Attr.td
955 ↗(On Diff #205503)

I think __hip_pinned_shadow__ best describes such variables. I will update the patch to adopt it.

yaxunl updated this revision to Diff 206504.Tue, Jun 25, 12:21 PM
yaxunl retitled this revision from [HIP] Support device_shadow variable to [HIP] Support attribute hip_pinned_shadow.
yaxunl edited the summary of this revision. (Show Details)

rename the attribute and make it HIP only.

tra accepted this revision.Tue, Jun 25, 1:31 PM

LGTM. Thank you!

This revision is now accepted and ready to land.Tue, Jun 25, 1:31 PM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptTue, Jun 25, 8:47 PM