[PATCH] D63020: [HIP] Fix visibility for 'extern' device variables.

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
hliao added a reviewer: yaxunl.
Herald added subscribers: cfe-commits, nhaehnle, jvesely.
Herald added a project: clang.

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


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D63020

Files:
  clang/lib/CodeGen/CodeGenModule.cpp
  clang/test/CodeGenCUDA/amdgpu-visibility.cu


Index: clang/test/CodeGenCUDA/amdgpu-visibility.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-visibility.cu
+++ clang/test/CodeGenCUDA/amdgpu-visibility.cu
@@ -13,6 +13,16 @@
 __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()
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3464,6 +3464,9 @@
     }
   }
 
+  if (GV->isDeclaration())
+    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
@@ -3473,9 +3476,6 @@
     return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace,
                                                        ExpectedAS, Ty);
 
-  if (GV->isDeclaration())
-    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
-
   return GV;
 }
 



Index: clang/test/CodeGenCUDA/amdgpu-visibility.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-visibility.cu
+++ clang/test/CodeGenCUDA/amdgpu-visibility.cu
@@ -13,6 +13,16 @@
 __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()
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3464,6 +3464,9 @@
     }
   }
 
+  if (GV->isDeclaration())
+    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
@@ -3473,9 +3476,6 @@
     return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace,
                                                        ExpectedAS, Ty);
 
-  if (GV->isDeclaration())
-    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
-
   return GV;
 }
 

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits