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

[Michael Liao via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rG0a220de9e9ca: [HIP] Fix visibility for 'extern' 
device variables. (authored by hliao).

Repository:
  rG LLVM Github Monorepo

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

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
@@ -3575,6 +3575,9 @@
     }
   }
 
+  if (GV->isDeclaration())
+    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
@@ -3584,9 +3587,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
@@ -3575,6 +3575,9 @@
     }
   }
 
+  if (GV->isDeclaration())
+    getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
@@ -3584,9 +3587,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