hliao created this revision.
hliao added reviewers: yaxunl, bogner.
Herald added subscribers: ormris, hiraditya.
hliao requested review of this revision.
Herald added projects: clang, LLVM.
Herald added subscribers: llvm-commits, cfe-commits.
- ``externally_initialized`` variables would be initialized or modified
elsewhere. Particularly, CUDA or HIP may have host code to initialize or modify
``externally_initialized`` device variables, which may not be explicitly
referenced on the device side but may still be used through the host side
interfaces. Not preserving them triggers the elimination of them in the
GlobalDCE and breaks the user code.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D105135
Files:
clang/test/CodeGenCUDA/host-used-device-var.cu
clang/test/CodeGenCUDA/unused-global-var.cu
llvm/lib/Transforms/IPO/Internalize.cpp
llvm/test/Transforms/Internalize/externally-initialized.ll
Index: llvm/test/Transforms/Internalize/externally-initialized.ll
===================================================================
--- /dev/null
+++ llvm/test/Transforms/Internalize/externally-initialized.ll
@@ -0,0 +1,7 @@
+; RUN: opt < %s -internalize -S | FileCheck %s
+; RUN: opt < %s -passes=internalize -S | FileCheck %s
+
+; CHECK: @G0
+; CHECK-NOT: internal
+; CHECK-SAME: global i32
+@G0 = protected externally_initialized global i32 0, align 4
Index: llvm/lib/Transforms/IPO/Internalize.cpp
===================================================================
--- llvm/lib/Transforms/IPO/Internalize.cpp
+++ llvm/lib/Transforms/IPO/Internalize.cpp
@@ -101,6 +101,12 @@
if (GV.hasDLLExportStorageClass())
return true;
+ // As the name suggestes, externally initialized variables need preserving as
+ // they would be initialized elsewhere externally.
+ if (const auto *G = dyn_cast<GlobalVariable>(&GV))
+ if (G->isExternallyInitialized())
+ return true;
+
// Already local, has nothing to do.
if (GV.hasLocalLinkage())
return false;
Index: clang/test/CodeGenCUDA/unused-global-var.cu
===================================================================
--- clang/test/CodeGenCUDA/unused-global-var.cu
+++ clang/test/CodeGenCUDA/unused-global-var.cu
@@ -17,12 +17,6 @@
// Check unused device/constant variables are eliminated.
-// NEGCHK-NOT: @v1
-__device__ int v1;
-
-// NEGCHK-NOT: @v2
-__constant__ int v2;
-
// NEGCHK-NOT: @_ZL2v3
constexpr int v3 = 1;
Index: clang/test/CodeGenCUDA/host-used-device-var.cu
===================================================================
--- clang/test/CodeGenCUDA/host-used-device-var.cu
+++ clang/test/CodeGenCUDA/host-used-device-var.cu
@@ -17,12 +17,6 @@
// Check device variables used by neither host nor device functioins are not
kept.
-// DEV-NEG-NOT: @v1
-__device__ int v1;
-
-// DEV-NEG-NOT: @v2
-__constant__ int v2;
-
// DEV-NEG-NOT: @_ZL2v3
static __device__ int v3;
Index: llvm/test/Transforms/Internalize/externally-initialized.ll
===================================================================
--- /dev/null
+++ llvm/test/Transforms/Internalize/externally-initialized.ll
@@ -0,0 +1,7 @@
+; RUN: opt < %s -internalize -S | FileCheck %s
+; RUN: opt < %s -passes=internalize -S | FileCheck %s
+
+; CHECK: @G0
+; CHECK-NOT: internal
+; CHECK-SAME: global i32
+@G0 = protected externally_initialized global i32 0, align 4
Index: llvm/lib/Transforms/IPO/Internalize.cpp
===================================================================
--- llvm/lib/Transforms/IPO/Internalize.cpp
+++ llvm/lib/Transforms/IPO/Internalize.cpp
@@ -101,6 +101,12 @@
if (GV.hasDLLExportStorageClass())
return true;
+ // As the name suggestes, externally initialized variables need preserving as
+ // they would be initialized elsewhere externally.
+ if (const auto *G = dyn_cast<GlobalVariable>(&GV))
+ if (G->isExternallyInitialized())
+ return true;
+
// Already local, has nothing to do.
if (GV.hasLocalLinkage())
return false;
Index: clang/test/CodeGenCUDA/unused-global-var.cu
===================================================================
--- clang/test/CodeGenCUDA/unused-global-var.cu
+++ clang/test/CodeGenCUDA/unused-global-var.cu
@@ -17,12 +17,6 @@
// Check unused device/constant variables are eliminated.
-// NEGCHK-NOT: @v1
-__device__ int v1;
-
-// NEGCHK-NOT: @v2
-__constant__ int v2;
-
// NEGCHK-NOT: @_ZL2v3
constexpr int v3 = 1;
Index: clang/test/CodeGenCUDA/host-used-device-var.cu
===================================================================
--- clang/test/CodeGenCUDA/host-used-device-var.cu
+++ clang/test/CodeGenCUDA/host-used-device-var.cu
@@ -17,12 +17,6 @@
// Check device variables used by neither host nor device functioins are not kept.
-// DEV-NEG-NOT: @v1
-__device__ int v1;
-
-// DEV-NEG-NOT: @v2
-__constant__ int v2;
-
// DEV-NEG-NOT: @_ZL2v3
static __device__ int v3;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
