[HIP] Fix visibility for 'extern' device variables.

[hliao]

Summary

• Fix a bug which misses the change for a variable to be set with target-specific attributes.

[hliao]

Sam, could you review this? Even though it has no functionality issue so far, from the code sequence, once there's an addrspacecast is inserted, we lose the chance to set target specific attributes if any.

[yaxunl]

LGTM. Thanks! @scchan This may fix the undefined symbol in work item struct issue at -O0

clang/lib/CodeGen/CodeGenModule.cpp

if (Context.getLangOpts().CPlusPlus && GV->hasExternalLinkage() &&
             }
             emitter.finalize(GV);
           }
         }
       }
     }
   }
 
+  if (GV->isDeclaration())
+    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
   assert(getContext().getTargetAddressSpace(ExpectedAS) ==
          Ty->getPointerAddressSpace());
   if (AddrSpace != ExpectedAS)
     return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace,
                                                        ExpectedAS, Ty);
 
-  if (GV->isDeclaration())
-    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
-
   return GV;
 }
 
 llvm::Constant *
 CodeGenModule::GetAddrOfGlobal(GlobalDecl GD,
                                ForDefinition_t IsForDefinition) {
   const Decl *D = GD.getDecl();
   if (isa<CXXConstructorDecl>(D) || isa<CXXDestructorDecl>(D))

clang/test/CodeGenCUDA/amdgpu-visibility.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility default -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-DEFAULT %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility protected -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-PROTECTED %s
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fapply-global-visibility-to-externs -fvisibility hidden -emit-llvm -o - %s | FileCheck --check-prefix=CHECK-HIDDEN %s
 
 #include "Inputs/cuda.h"
 
 // CHECK-DEFAULT: @c = addrspace(4) externally_initialized global
 // CHECK-DEFAULT: @g = addrspace(1) externally_initialized global
 // CHECK-PROTECTED: @c = protected addrspace(4) externally_initialized global
 // CHECK-PROTECTED: @g = protected addrspace(1) externally_initialized global
 // CHECK-HIDDEN: @c = protected addrspace(4) externally_initialized global
 // CHECK-HIDDEN: @g = protected addrspace(1) externally_initialized global
 __constant__ int c;
 __device__ int g;
 
+// CHECK-DEFAULT: @e = external addrspace(1) global
+// CHECK-PROTECTED: @e = external protected addrspace(1) global
+// CHECK-HIDDEN: @e = external protected addrspace(1) global
+extern __device__ int e;
+
+// dummy one to hold reference to `e`.
+__device__ int f() {
+  return e;
+}
+
 // CHECK-DEFAULT: define amdgpu_kernel void @_Z3foov()
 // CHECK-PROTECTED: define protected amdgpu_kernel void @_Z3foov()
 // CHECK-HIDDEN: define protected amdgpu_kernel void @_Z3foov()
 __global__ void foo() {
   g = c;
 }