[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 ↗

(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]

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 ↗

(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]

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.

[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

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]

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]

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]

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.

cfe/trunk/include/clang/Basic/Attr.td

class LangOpt<string name, code customCode = [{}]> {
 
   // A custom predicate, written as an expression evaluated in a context with
   // "LangOpts" bound.
   code CustomCode = customCode;
 }
 def MicrosoftExt : LangOpt<"MicrosoftExt">;
 def Borland : LangOpt<"Borland">;
 def CUDA : LangOpt<"CUDA">;
+def HIP : LangOpt<"HIP">;
 def COnly : LangOpt<"COnly", "!LangOpts.CPlusPlus">;
 def CPlusPlus : LangOpt<"CPlusPlus">;
 def OpenCL : LangOpt<"OpenCL">;
 def RenderScript : LangOpt<"RenderScript">;
 def ObjC : LangOpt<"ObjC">;
 def BlocksSupported : LangOpt<"Blocks">;
 def ObjCAutoRefCount : LangOpt<"ObjCAutoRefCount">;
 def ObjCNonFragileRuntime : LangOpt<"ObjCNonFragileRuntime",

 
 def CUDADevice : InheritableAttr {
   let Spellings = [GNU<"device">, Declspec<"__device__">];
   let Subjects = SubjectList<[Function, Var]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
+def HIPPinnedShadow : InheritableAttr {
+  let Spellings = [GNU<"hip_pinned_shadow">, Declspec<"__hip_pinned_shadow__">];
+  let Subjects = SubjectList<[Var]>;
+  let LangOpts = [HIP];
+  let Documentation = [HIPPinnedShadowDocs];
+}
+
 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__">];

cfe/trunk/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 HIPPinnedShadowDocs : Documentation {
+  let Category = DocCatType;
+  let Content = [{
+The GNU style attribute __attribute__((hip_pinned_shadow)) or MSVC style attribute
+__declspec(hip_pinned_shadow) can be added to the definition of a global variable
+to indicate it is a HIP pinned shadow variable. A HIP pinned shadow variable can
+be accessed on both device side and host side. It has external linkage and is
+not initialized on device side. It has internal linkage and is initialized by
+the initializer on host side.
+  }];
+}
 No newline at end of file

cfe/trunk/lib/CodeGen/CodeGenModule.cpp

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 && Global->hasAttr<HIPPinnedShadowAttr>()))
         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 pinned shadow of initialized host-side global variables are also
+  // left undefined.
+  bool IsHIPPinnedShadowVar =
+      getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
+  if (getLangOpts().CUDA &&
+      (IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
     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>() ||
+          D->hasAttr<HIPPinnedShadowAttr>()) {
         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>())
           Flags |= CGCUDARuntime::ConstantDeviceVar;
         // Extern global variables will be registered in the TU where they are
         // defined.
         if (!D->hasExternalStorage())
           getCUDARuntime().registerDeviceVar(D, *GV, Flags);
       } else if (D->hasAttr<CUDASharedAttr>())
         // __shared__ variables are odd. Shadows do get created, but
         // they are not registered with the CUDA runtime, so they
         // can't really be used to access their device-side
         // counterparts. It's not clear yet whether it's nvcc's bug or
         // a feature, but we've got to do the same for compatibility.
         Linkage = llvm::GlobalValue::InternalLinkage;
     }
   }
 
+  if (!IsHIPPinnedShadowVar)
     GV->setInitializer(Init);
   if (emitter) emitter->finalize(GV);
 
   // If it is safe to mark the global 'constant', do so now.
   GV->setConstant(!NeedsGlobalCtor && !NeedsGlobalDtor &&
                   isTypeConstant(D->getType(), true));
 
   // If it is in a read-only section, mark it 'constant'.
   if (const SectionAttr *SA = D->getAttr<SectionAttr>()) {

cfe/trunk/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>() ||
+           D->hasAttr<HIPPinnedShadowAttr>()));
+}
+
+static bool requiresAMDGPUDefaultVisibility(const Decl *D,
+                                            llvm::GlobalValue *GV) {
+  if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
+    return false;
+
+  return isa<VarDecl>(D) && D->hasAttr<HIPPinnedShadowAttr>();
 }
 
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(
     const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
-  if (requiresAMDGPUProtectedVisibility(D, GV)) {
+  if (requiresAMDGPUDefaultVisibility(D, GV)) {
+    GV->setVisibility(llvm::GlobalValue::DefaultVisibility);
+    GV->setDSOLocal(false);
+  } else if (requiresAMDGPUProtectedVisibility(D, GV)) {
     GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
     GV->setDSOLocal(true);
   }
 
   if (GV->isDeclaration())
     return;
   const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D);
   if (!FD)

cfe/trunk/lib/Driver/ToolChains/HIP.cpp

 
 void AMDGCN::Linker::constructLldCommand(Compilation &C, const JobAction &JA,
                                           const InputInfoList &Inputs,
                                           const InputInfo &Output,
                                           const llvm::opt::ArgList &Args,
                                           const char *InputFileName) const {
   // Construct lld command.
   // The output from ld.lld is an HSA code object file.
-  ArgStringList LldArgs{"-flavor",    "gnu", "--no-undefined",
-                        "-shared",    "-o",  Output.getFilename(),
-                        InputFileName};
+  ArgStringList LldArgs{
+      "-flavor", "gnu", "-shared", "-o", Output.getFilename(), InputFileName};
   SmallString<128> LldPath(C.getDriver().Dir);
   llvm::sys::path::append(LldPath, "lld");
   const char *Lld = Args.MakeArgString(LldPath);
   C.addCommand(llvm::make_unique<Command>(JA, *this, Lld, LldArgs, Inputs));
 }
 
 // Construct a clang-offload-bundler command to bundle code objects for
 // different GPU's into a HIP fat binary.

cfe/trunk/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_HIPPinnedShadow:
+    handleSimpleAttributeWithExclusions<HIPPinnedShadowAttr, CUDADeviceAttr,
+                                        CUDAConstantAttr>(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);

cfe/trunk/test/AST/ast-dump-hip-pinned-shadow.cu

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

cfe/trunk/test/CodeGenCUDA/hip-pinned-shadow.cu

+// REQUIRES: amdgpu-registered-target
+
+// 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 texture : public textureReference {
+texture() { a = 1; }
+};
+
+__attribute__((hip_pinned_shadow)) texture<float, 2, 1> tex;
+// CUDADEV-NOT: @tex
+// CUDAHOST-NOT: call i32 @__hipRegisterVar{{.*}}@tex
+// HIPDEV: @tex = external addrspace(1) global %struct.texture
+// HIPDEV-NOT: declare{{.*}}void @_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST:  define{{.*}}@_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST:  call i32 @__hipRegisterVar{{.*}}@tex{{.*}}i32 0, i32 4, i32 0, i32 0)

cfe/trunk/test/Driver/hip-toolchain-no-rdc.hip

 // CHECK: [[OPT:".*opt"]] [[LINKED_BC_DEV_A_803]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-mcpu=gfx803"
 // CHECK-SAME: "-o" [[OPT_BC_DEV_A_803:".*-gfx803-optimized.*bc"]]
 
 // CHECK: [[LLC: ".*llc"]] [[OPT_BC_DEV_A_803]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-filetype=obj"
 // CHECK-SAME: "-mcpu=gfx803" "-o" [[OBJ_DEV_A_803:".*-gfx803-.*o"]]
 
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
 // CHECK-SAME: "-o" "[[IMG_DEV_A_803:.*out]]" [[OBJ_DEV_A_803]]
 
 //
 // Compile device code in a.cu to code object for gfx900.
 //
 
 // CHECK: [[CLANG:".*clang.*"]] "-cc1" "-triple" "amdgcn-amd-amdhsa"
 // CHECK-SAME: "-aux-triple" "x86_64-unknown-linux-gnu"

 // CHECK: [[OPT:".*opt"]] [[LINKED_BC_DEV_A_900]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-mcpu=gfx900"
 // CHECK-SAME: "-o" [[OPT_BC_DEV_A_900:".*-gfx900-optimized.*bc"]]
 
 // CHECK: [[LLC: ".*llc"]] [[OPT_BC_DEV_A_900]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-filetype=obj"
 // CHECK-SAME: "-mcpu=gfx900" "-o" [[OBJ_DEV_A_900:".*-gfx900-.*o"]]
 
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
 // CHECK-SAME: "-o" "[[IMG_DEV_A_900:.*out]]" [[OBJ_DEV_A_900]]
 
 //
 // Bundle and embed device code in host object for a.cu.
 //
 
 // CHECK: [[BUNDLER:".*clang-offload-bundler"]] "-type=o"
 // CHECK-SAME: "-targets={{.*}},hip-amdgcn-amd-amdhsa-gfx803,hip-amdgcn-amd-amdhsa-gfx900"

 // CHECK: [[OPT:".*opt"]] [[LINKED_BC_DEV_B_803]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-mcpu=gfx803"
 // CHECK-SAME: "-o" [[OPT_BC_DEV_B_803:".*-gfx803-optimized.*bc"]]
 
 // CHECK: [[LLC: ".*llc"]] [[OPT_BC_DEV_B_803]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-filetype=obj"
 // CHECK-SAME: "-mcpu=gfx803" "-o" [[OBJ_DEV_B_803:".*-gfx803-.*o"]]
 
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
 // CHECK-SAME: "-o" "[[IMG_DEV_B_803:.*out]]" [[OBJ_DEV_B_803]]
 
 //
 // Compile device code in b.hip to code object for gfx900.
 //
 
 // CHECK: [[CLANG:".*clang.*"]] "-cc1" "-triple" "amdgcn-amd-amdhsa"
 // CHECK-SAME: "-aux-triple" "x86_64-unknown-linux-gnu"

 // CHECK: [[OPT:".*opt"]] [[LINKED_BC_DEV_B_900]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-mcpu=gfx900"
 // CHECK-SAME: "-o" [[OPT_BC_DEV_B_900:".*-gfx900-optimized.*bc"]]
 
 // CHECK: [[LLC: ".*llc"]] [[OPT_BC_DEV_B_900]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-filetype=obj"
 // CHECK-SAME: "-mcpu=gfx900" "-o" [[OBJ_DEV_B_900:".*-gfx900-.*o"]]
 
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
 // CHECK-SAME: "-o" "[[IMG_DEV_B_900:.*out]]" [[OBJ_DEV_B_900]]
 
 //
 // Bundle and embed device code in host object for b.hip.
 //
 
 // CHECK: [[BUNDLER:".*clang-offload-bundler"]] "-type=o"
 // CHECK-SAME: "-targets={{.*}},hip-amdgcn-amd-amdhsa-gfx803,hip-amdgcn-amd-amdhsa-gfx900"

cfe/trunk/test/Driver/hip-toolchain-rdc.hip

 // CHECK: [[OPT:".*opt"]] [[LINKED_BC_DEV1]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-mcpu=gfx803"
 // CHECK-SAME: "-o" [[OPT_BC_DEV1:".*-gfx803-optimized.*bc"]]
 
 // CHECK: [[LLC: ".*llc"]] [[OPT_BC_DEV1]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-filetype=obj"
 // CHECK-SAME: "-mcpu=gfx803" "-o" [[OBJ_DEV1:".*-gfx803-.*o"]]
 
-// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD: ".*lld"]] "-flavor" "gnu" "-shared"
 // CHECK-SAME: "-o" "[[IMG_DEV1:.*out]]" [[OBJ_DEV1]]
 
 // CHECK: [[CLANG]] "-cc1" "-triple" "amdgcn-amd-amdhsa" 
 // CHECK-SAME: "-aux-triple" "x86_64-unknown-linux-gnu"
 // CHECK-SAME: "-emit-llvm-bc"
 // CHECK-SAME: {{.*}} "-main-file-name" "a.cu" {{.*}} "-target-cpu" "gfx900"
 // CHECK-SAME: "-fcuda-is-device" "-fgpu-rdc"
 // CHECK-SAME: "{{.*}}lib1.bc" "{{.*}}lib2.bc"

 // CHECK: [[OPT]] [[LINKED_BC_DEV2]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-mcpu=gfx900"
 // CHECK-SAME: "-o" [[OPT_BC_DEV2:".*-gfx900-optimized.*bc"]]
 
 // CHECK: [[LLC]] [[OPT_BC_DEV2]] "-mtriple=amdgcn-amd-amdhsa"
 // CHECK-SAME: "-filetype=obj"
 // CHECK-SAME: "-mcpu=gfx900" "-o" [[OBJ_DEV2:".*-gfx900-.*o"]]
 
-// CHECK: [[LLD]] "-flavor" "gnu" "--no-undefined" "-shared"
+// CHECK: [[LLD]] "-flavor" "gnu" "-shared"
 // CHECK-SAME: "-o" "[[IMG_DEV2:.*out]]" [[OBJ_DEV2]]
 
 // CHECK: [[CLANG]] "-cc1" "-triple" "x86_64-unknown-linux-gnu"
 // CHECK-SAME: "-aux-triple" "amdgcn-amd-amdhsa"
 // CHECK-SAME: "-emit-obj"
 // CHECK-SAME: {{.*}} "-main-file-name" "a.cu"
 // CHECK-SAME: {{.*}} "-o" [[A_OBJ_HOST:".*o"]] "-x" "hip"
 // CHECK-SAME: {{.*}} [[A_SRC]]

cfe/trunk/test/Misc/pragma-attribute-supported-attributes-list.test

 // CHECK-NEXT: DisableTailCalls (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: EnableIf (SubjectMatchRule_function)
 // CHECK-NEXT: EnumExtensibility (SubjectMatchRule_enum)
 // CHECK-NEXT: ExcludeFromExplicitInstantiation (SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_record)
 // CHECK-NEXT: ExternalSourceSymbol ((SubjectMatchRule_record, SubjectMatchRule_enum, SubjectMatchRule_enum_constant, SubjectMatchRule_field, SubjectMatchRule_function, SubjectMatchRule_namespace, SubjectMatchRule_objc_category, SubjectMatchRule_objc_implementation, SubjectMatchRule_objc_interface, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property, SubjectMatchRule_objc_protocol, SubjectMatchRule_record, SubjectMatchRule_type_alias, SubjectMatchRule_variable))
 // CHECK-NEXT: FlagEnum (SubjectMatchRule_enum)
 // CHECK-NEXT: Flatten (SubjectMatchRule_function)
 // CHECK-NEXT: GNUInline (SubjectMatchRule_function)
+// CHECK-NEXT: HIPPinnedShadow (SubjectMatchRule_variable)
 // CHECK-NEXT: Hot (SubjectMatchRule_function)
 // CHECK-NEXT: IBAction (SubjectMatchRule_objc_method_is_instance)
 // CHECK-NEXT: IFunc (SubjectMatchRule_function)
 // CHECK-NEXT: InitPriority (SubjectMatchRule_variable)
 // CHECK-NEXT: InternalLinkage (SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_record)
 // CHECK-NEXT: LTOVisibilityPublic (SubjectMatchRule_record)
 // CHECK-NEXT: Lockable (SubjectMatchRule_record)
 // CHECK-NEXT: MIGServerRoutine (SubjectMatchRule_function, SubjectMatchRule_objc_method, SubjectMatchRule_block)

cfe/trunk/test/SemaCUDA/hip-pinned-shadow.cu

+// 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 -fsyntax-only -verify
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN:     -emit-llvm -o - -x hip %s -fsyntax-only -verify
+
+#define __device__ __attribute__((device))
+#define __constant__ __attribute__((constant))
+#define __hip_pinned_shadow__ __attribute((hip_pinned_shadow))
+
+struct textureReference {
+  int a;
+};
+
+template <class T, int texType, int hipTextureReadMode>
+struct texture : public textureReference {
+texture() { a = 1; }
+};
+
+__hip_pinned_shadow__ texture<float, 2, 1> tex;
+__device__ __hip_pinned_shadow__ texture<float, 2, 1> tex2; // expected-error{{'hip_pinned_shadow' and 'device' attributes are not compatible}}
+                                                            // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}}
+                                                            // expected-note@-2{{conflicting attribute is here}}
+__constant__ __hip_pinned_shadow__ texture<float, 2, 1> tex3; // expected-error{{'hip_pinned_shadow' and 'constant' attributes are not compatible}}
+                                                              // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables}}
+                                                              // expected-note@-2{{conflicting attribute is here}}