[HIP] Support attribute hip_pinned_shadow

[yaxunl]

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.

[yaxunl]

ping

[tra]

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

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

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

3781

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]

Revised by Artem's comments.

[tra]

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?

[yaxunl]

Revised by Artem's comments.

[yaxunl]

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]

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

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

3781

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]

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

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.

[yaxunl]

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]

lib/CodeGen/CodeGenModule.cpp

3781

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]

include/clang/Basic/Attr.td

954

HIPDeviceShadow ?

955

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

Shis should probably be [HIP] now, too.

include/clang/Basic/AttrDocs.td

4164–4171

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

3781

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]

include/clang/Basic/Attr.td

955

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

[yaxunl]

rename the attribute and make it HIP only.

[tra]

LGTM. Thank you!

[yaxunl]

@tra Originally we introduced the __hip_pinned_shadow__ attribute to support texture type. It ended up as a variable attribute for more generic usage.

However we got usability issue since users have to add this attribute whenever they use texture type. I would like to propose to allow this attribute to both types and variables.

What do you think?

Thanks.

include/clang/Basic/Attr.td

 
 def CUDADevice : InheritableAttr {
   let Spellings = [GNU<"device">, Declspec<"__device__">];
   let Subjects = SubjectList<[Function, Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDADeviceBuiltin : IgnoredAttr {
   let Spellings = [GNU<"device_builtin">, Declspec<"__device_builtin__">];
   let LangOpts = [CUDA];
 }
 
 def CUDADeviceBuiltinSurfaceType : IgnoredAttr {
   let Spellings = [GNU<"device_builtin_surface_type">,
                    Declspec<"__device_builtin_surface_type__">];
   let LangOpts = [CUDA];
 }
 
-def CUDADeviceBuiltinTextureType : IgnoredAttr {
+def CUDADeviceBuiltinTextureType : TypeAttr {
   let Spellings = [GNU<"device_builtin_texture_type">,
                    Declspec<"__device_builtin_texture_type__">];
   let LangOpts = [CUDA];
+  let Documentation = [DeviceBuiltinTextureTypeDocs];
 }
 
 def CUDAGlobal : InheritableAttr {
   let Spellings = [GNU<"global">, Declspec<"__global__">];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }

include/clang/Basic/AttrDocs.td

 Studio debugger uses this information to `profile memory usage`_.
 
 .. _profile memory usage: https://docs.microsoft.com/en-us/visualstudio/profiling/memory-usage
 
 This attribute does not affect optimizations in any way, unlike GCC's
 ``__attribute__((malloc))``.
 }];
 }
+
+def DeviceBuiltinTextureTypeDocs : Documentation {
+  let Category = DocCatType;
+  let Content = [{
+The GNU style attribute __attribute__((device_builtin_texture_type)) or MSVC
+style attribute __declspec(device_builtin_texture_type) can be added to the
+definition of a class template to indicate it is the HIP texture reference type,
+which is defined as
+
+   .. code-block:: c++
+
+  template <class T, int texType, enum hipTextureReadMode>
+    struct __attribute__((device_builtin_texture_type)) texture
+      : public textureReference { ... }
+
+where T is the data type of texels, texType is the type of the texture, e.g 1D,
+2D, or 3D, hipTextureReadMode is the enumeration type of the read mode of
+the texture. HIP texture reference type is defined as a class template so that
+the template arguments are compile time constants. HIP texture reference also
+contains information which can be changed at run time, e.g. filter mode, address
+mode, which are defined in the base class textureReference. HIP texture
+references can only be defined as global variables. In device code, they can
+be passed to texture fetch API functions to fetch texels. In host code, they
+can be modified and binded to textures by texture binding API functions. They
+should be binded in host code before they can be used in kernels. An example
+of HIP texture reference is as follows:
+
+   .. code-block:: c++
+   texture<float, 2, hipReadModeElementType> tex;
+
+   __global__ void tex2DKernel(float* outputData,
+                               int width, int height) {
+     int x = blockIdx.x * blockDim.x + threadIdx.x;
+     int y = blockIdx.y * blockDim.y + threadIdx.y;
+     outputData[y * width + x] = tex2D(tex, x, y);
+   }
+
+   int main() {
+     // ...
+     tex.addressMode[0] = hipAddressModeWrap;
+     tex.addressMode[1] = hipAddressModeWrap;
+     tex.filterMode = hipFilterModePoint;
+     tex.normalized = 0;
+     hipBindTextureToArray(tex, hipArray, channelDesc);
+     tex2DKernel<<<dim3(dimGrid), dim3(dimBlock), 0, 0>>>(dData, width, height);
+     // ...
+   }
+}
+
+It is ignored for CUDA and other languages.
+  }];
+}
 No newline at end of file

lib/AST/TypePrinter.cpp

#include "clang/Basic/AttrList.inc"
   case attr::ObjCInertUnsafeUnretained:
   case attr::ObjCKindOf:
   case attr::ObjCOwnership:
   case attr::Ptr32:
   case attr::Ptr64:
   case attr::SPtr:
   case attr::UPtr:
   case attr::AddressSpace:
+  case attr::CUDADeviceBuiltinTextureType:
     llvm_unreachable("This attribute should have been handled already");
 
   case attr::NSReturnsRetained:
     OS << "ns_returns_retained";
     break;
 
   // FIXME: When Sema learns to form this AttributedType, avoid printing the
   // attribute again in printFunctionProtoAfter.

lib/CodeGen/CodeGenModule.cpp

ConstantAddress CodeGenModule::GetWeakRefReference(const ValueDecl *VD) {
 
   auto *F = cast<llvm::GlobalValue>(Aliasee);
   F->setLinkage(llvm::Function::ExternalWeakLinkage);
   WeakRefReferences.insert(F);
 
   return ConstantAddress(Aliasee, Alignment);
 }
 
+static bool isCUDATextureType(const QualType &T) {
+  return T->isRecordType() && T->getAs<RecordType>()
+                                  ->getDecl()
+                                  ->hasAttr<CUDADeviceBuiltinTextureTypeAttr>();
+}
+
 void CodeGenModule::EmitGlobal(GlobalDecl GD) {
   const auto *Global = cast<ValueDecl>(GD.getDecl());
 
   // Weak references don't produce any output by themselves.
   if (Global->hasAttr<WeakRefAttr>())
     return;
 
   // If this is an alias definition (which otherwise looks like a declaration)
if (Global->hasAttr<CPUDispatchAttr>())
     return emitCPUDispatchDefinition(GD);
 
   // If this is CUDA, be selective about which declarations we emit.
   if (LangOpts.CUDA) {
     if (LangOpts.CUDAIsDevice) {
       if (!Global->hasAttr<CUDADeviceAttr>() &&
           !Global->hasAttr<CUDAGlobalAttr>() &&
           !Global->hasAttr<CUDAConstantAttr>() &&
-          !Global->hasAttr<CUDASharedAttr>())
+          !Global->hasAttr<CUDASharedAttr>() &&
+          !(LangOpts.HIP && isCUDATextureType(Global->getType())))
         return;
     } else {
       // We need to emit host-side 'shadows' for all global
       // device-side variables because the CUDA runtime needs their
       // size and host-side address in order to provide access to
       // their device-side incarnations.
 
       // So device-only functions are the only things we skip.
void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   bool IsCUDASharedVar =
       getLangOpts().CUDAIsDevice && D->hasAttr<CUDASharedAttr>();
   // Shadows of initialized device-side global variables are also left
   // undefined.
   bool IsCUDAShadowVar =
       !getLangOpts().CUDAIsDevice &&
       (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
        D->hasAttr<CUDASharedAttr>());
-  if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar))
+  // HIP texture references have non-trivial default constructors, therefore
+  // they cannot be initialized in device code.
+  if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar ||
+                             (getLangOpts().HIP && getLangOpts().CUDAIsDevice &&
+                              isCUDATextureType(D->getType()))))
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   else if (!InitExpr) {
     // This is a tentative definition; tentative definitions are
     // implicitly initialized with { 0 }.
     //
     // Note that tentative definitions are only emitted at the end of
     // a translation unit, so they should never have incomplete
     // type. In addition, EmitTentativeDefinition makes sure that we
if (LangOpts.CUDAIsDevice) {
       if (Linkage != llvm::GlobalValue::InternalLinkage &&
           (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
         GV->setExternallyInitialized(true);
     } else {
       // Host-side shadows of external declarations of device-side
       // global variables become internal definitions. These have to
       // be internal in order to prevent name conflicts with global
       // host variables with the same name in a different TUs.
-      if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
+      if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+          (isCUDATextureType(D->getType()) && LangOpts.HIP)) {
         Linkage = llvm::GlobalValue::InternalLinkage;
 
         // Shadow variables and their properties must be registered
         // with CUDA runtime.
         unsigned Flags = 0;
         if (!D->hasDefinition())
           Flags |= CGCUDARuntime::ExternDeviceVar;
         if (D->hasAttr<CUDAConstantAttr>())

lib/CodeGen/TargetInfo.cpp

 static bool requiresAMDGPUProtectedVisibility(const Decl *D,
                                               llvm::GlobalValue *GV) {
   if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
     return false;
 
   return D->hasAttr<OpenCLKernelAttr>() ||
          (isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
          (isa<VarDecl>(D) &&
-          (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()));
+          (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+           (cast<VarDecl>(D)->getType()->isRecordType() &&
+            cast<VarDecl>(D)
+                ->getType()
+                ->getAs<RecordType>()
+                ->getDecl()
+                ->hasAttr<CUDADeviceBuiltinTextureTypeAttr>())));
 }
 
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(
     const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
   if (requiresAMDGPUProtectedVisibility(D, GV)) {
     GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
     GV->setDSOLocal(true);
   }

lib/Sema/SemaDeclAttr.cpp

case ParsedAttr::AT_CUDAGlobal:
     break;
   case ParsedAttr::AT_CUDADevice:
     handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D,
                                                                         AL);
     break;
   case ParsedAttr::AT_CUDAHost:
     handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_CUDADeviceBuiltinTextureType:
+    handleSimpleAttribute<CUDADeviceBuiltinTextureTypeAttr>(S, D, AL);
+    break;
   case ParsedAttr::AT_GNUInline:
     handleGNUInlineAttr(S, D, AL);
     break;
   case ParsedAttr::AT_CUDALaunchBounds:
     handleLaunchBoundsAttr(S, D, AL);
     break;
   case ParsedAttr::AT_Restrict:
     handleRestrictAttr(S, D, AL);

test/AST/ast-dump-cuda-texture.cu

+// RUN: %clang_cc1 -fcuda-is-device -ast-dump -ast-dump-filter texture %s | FileCheck -strict-whitespace %s
+// RUN: %clang_cc1 -ast-dump -ast-dump-filter texture %s | FileCheck -strict-whitespace %s
+struct textureReference {
+  int a;
+};
+
+// CHECK: CUDADeviceBuiltinTextureTypeAttr
+template <class T, int texType, int hipTextureReadMode>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+texture() { a = 1; }
+};

test/CodeGenCUDA/texture.cu

+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 -fcuda-is-device \
+// RUN:     -emit-llvm -o - %s | FileCheck -check-prefixes=CUDADEV %s
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN:     -emit-llvm -o - %s | FileCheck -check-prefixes=CUDAHOST %s
+
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 -fvisibility hidden -fapply-global-visibility-to-externs \
+// RUN:     -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPDEV %s
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN:     -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPHOST %s
+
+struct textureReference {
+  int a;
+};
+
+template <class T, int texType, int hipTextureReadMode>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+texture() { a = 1; }
+};
+
+texture<float, 2, 1> tex;
+// CUDADEV-NOT: @tex
+// CUDAHOST-NOT: call i32 @__hipRegisterVar{{.*}}@tex
+// HIPDEV: @tex = protected addrspace(1) global %struct.texture undef
+// HIPDEV-NOT: declare{{.*}}void @_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST:  define{{.*}}@_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST:  call i32 @__hipRegisterVar{{.*}}@tex

test/SemaCUDA/attr-declspec.cu

 // Test the __declspec spellings of CUDA attributes.
 //
 // RUN: %clang_cc1 -fsyntax-only -fms-extensions -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fms-extensions -fcuda-is-device -verify %s
 // Now pretend that we're compiling a C file. There should be warnings.
 // RUN: %clang_cc1 -DEXPECT_WARNINGS -fms-extensions -fsyntax-only -verify -x c %s
 
 #if defined(EXPECT_WARNINGS)
-// expected-warning@+12 {{'__device__' attribute ignored}}
-// expected-warning@+12 {{'__global__' attribute ignored}}
-// expected-warning@+12 {{'__constant__' attribute ignored}}
-// expected-warning@+12 {{'__shared__' attribute ignored}}
-// expected-warning@+12 {{'__host__' attribute ignored}}
+// expected-warning@+13 {{'__device__' attribute ignored}}
+// expected-warning@+13 {{'__global__' attribute ignored}}
+// expected-warning@+13 {{'__constant__' attribute ignored}}
+// expected-warning@+13 {{'__shared__' attribute ignored}}
+// expected-warning@+13 {{'__host__' attribute ignored}}
+// expected-warning@+19 {{'__device_builtin_texture_type__' attribute ignored}}
 //
 // (Currently we don't for the other attributes. They are implemented with
 // IgnoredAttr, which is ignored irrespective of any LangOpts.)
 #else
 // expected-no-diagnostics
 #endif
 
 __declspec(__device__) void f_device();

test/SemaCUDA/attributes-on-non-cuda.cu

 // Tests that CUDA attributes are warnings when compiling C files, but not when
 // compiling CUDA files.
 //
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
 // Now pretend that we're compiling a C file. There should be warnings.
 // RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s
 
 #if defined(EXPECT_WARNINGS)
-// expected-warning@+12 {{'device' attribute ignored}}
-// expected-warning@+12 {{'global' attribute ignored}}
-// expected-warning@+12 {{'constant' attribute ignored}}
-// expected-warning@+12 {{'shared' attribute ignored}}
-// expected-warning@+12 {{'host' attribute ignored}}
+// expected-warning@+13 {{'device' attribute ignored}}
+// expected-warning@+13 {{'global' attribute ignored}}
+// expected-warning@+13 {{'constant' attribute ignored}}
+// expected-warning@+13 {{'shared' attribute ignored}}
+// expected-warning@+13 {{'host' attribute ignored}}
+// expected-warning@+20 {{'device_builtin_texture_type' attribute ignored}}
 //
 // NOTE: IgnoredAttr in clang which is used for the rest of
 // attributes ignores LangOpts, so there are no warnings.
 #else
 // expected-no-diagnostics
 #endif
 
 __attribute__((device)) void f_device();