hliao created this revision.
hliao added a reviewer: yaxunl.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
- By declaring device variables as `static`, we assume they won't be
addressable from the host side. Thus, no `externally_initialized` is required.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D62603
Files:
clang/lib/CodeGen/CodeGenModule.cpp
clang/test/CodeGenCUDA/device-var-init.cu
Index: clang/test/CodeGenCUDA/device-var-init.cu
===================================================================
--- clang/test/CodeGenCUDA/device-var-init.cu
+++ clang/test/CodeGenCUDA/device-var-init.cu
@@ -33,6 +33,16 @@
// DEVICE: @d_v_i = addrspace(1) externally_initialized global i32 1,
// HOST: @d_v_i = internal global i32 undef,
+// For `static` device variables, assume they won't be addressed from the host
+// side.
+static __device__ int d_s_v_i = 1;
+// DEVICE: @_ZL7d_s_v_i = internal addrspace(1) global i32 1,
+
+// Dummy function to keep static variables referenced.
+__device__ int foo() {
+ return d_s_v_i;
+}
+
// trivial constructor -- allowed
__device__ T d_t;
// DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T
zeroinitializer
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3869,7 +3869,8 @@
// / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())."
if (GV && LangOpts.CUDA) {
if (LangOpts.CUDAIsDevice) {
- if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>())
+ if (Linkage != llvm::GlobalValue::InternalLinkage &&
+ (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
GV->setExternallyInitialized(true);
} else {
// Host-side shadows of external declarations of device-side
Index: clang/test/CodeGenCUDA/device-var-init.cu
===================================================================
--- clang/test/CodeGenCUDA/device-var-init.cu
+++ clang/test/CodeGenCUDA/device-var-init.cu
@@ -33,6 +33,16 @@
// DEVICE: @d_v_i = addrspace(1) externally_initialized global i32 1,
// HOST: @d_v_i = internal global i32 undef,
+// For `static` device variables, assume they won't be addressed from the host
+// side.
+static __device__ int d_s_v_i = 1;
+// DEVICE: @_ZL7d_s_v_i = internal addrspace(1) global i32 1,
+
+// Dummy function to keep static variables referenced.
+__device__ int foo() {
+ return d_s_v_i;
+}
+
// trivial constructor -- allowed
__device__ T d_t;
// DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -3869,7 +3869,8 @@
// / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())."
if (GV && LangOpts.CUDA) {
if (LangOpts.CUDAIsDevice) {
- if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>())
+ if (Linkage != llvm::GlobalValue::InternalLinkage &&
+ (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
GV->setExternallyInitialized(true);
} else {
// Host-side shadows of external declarations of device-side
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
