[PATCH] D62738: [HIP] Support texture type

[Artem Belevich via Phabricator via cfe-commits]

tra added a comment.

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?


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D62738/new/

https://reviews.llvm.org/D62738