[Internalize] Preserve variables externally initialized.

[hliao]

Summary

• `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.

[hliao]

clang/test/CodeGenCUDA/host-used-device-var.cu

20

BTW, as clang codegen tests, those checks should not rely on middle-end optimizations to work correctly.

clang/test/CodeGenCUDA/unused-global-var.cu

20

BTW, as clang codegen tests, those checks should not rely on middle-end optimizations to work correctly.

[hliao]

Fix typo. Kindly PING for review.

[tra]

clang/test/CodeGenCUDA/host-used-device-var.cu

20

What happens if you remove -O3?
Sometimes optimization is helpful as it reduces the amount of IR we need to check and how it's structured. It's much easier to see what went wrong when you get just a handful of IR statements in the failing test output.

21–24

These should be changed to positive checks to verify that they are emitted.
Ditto for other tests.

[hliao]

Kindly PING for review.

[yaxunl]

LGTM. Pls change the negative tests to positive tests as Artem suggested.

clang/test/CodeGenCUDA/host-used-device-var.cu

20

These tests were introduced with -O3 since they can only be tested with clang and -O3. (https://reviews.llvm.org/D98783) Clang may emit IR containing unused LLVM constants. Saving the IR as a file will cause such unused LLVM constants to disappear, therefore unable to reproduce the issue.

20

If -O3 is removed, those unused variables will not be internalized and removed.

[hliao]

Revert part of tests back and convert them into positive ones.

[yaxunl]

LGTM. Thanks.

clang/test/CodeGenCUDA/host-used-device-var.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
 // RUN:   -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
 // RUN:   | FileCheck -check-prefix=DEV-NEG %s
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
 // RUN:   -std=c++17 -O3 -emit-llvm -o - | FileCheck -check-prefix=HOST-NEG %s
 
 #include "Inputs/cuda.h"
 
-// Check device variables used by neither host nor device functioins are not kept.
-
-// DEV-NEG-NOT: @v1
+// DEV-DAG: @v1
 __device__ int v1;
 
-// DEV-NEG-NOT: @v2
+// DEV-DAG: @v2
 __constant__ int v2;
 
+// Check device variables used by neither host nor device functioins are not kept.
+
 // DEV-NEG-NOT: @_ZL2v3
 static __device__ int v3;
 
 // Check device variables used by host functions are kept.
 
 // DEV-DAG: @u1
 __device__ int u1;
 

clang/test/CodeGenCUDA/unused-global-var.cu

 
 // AMDGPU internalize unused global variables for whole-program compilation
 // (-fno-gpu-rdc for each TU, or -fgpu-rdc for LTO), which are then
 // eliminated by global DCE. If there are invisible unused address space casts
 // for global variables, these dead users need to be eliminated by global
 // DCE before internalization. This test makes sure unused global variables
 // are eliminated.
 
-// Check unused device/constant variables are eliminated.
-
-// NEGCHK-NOT: @v1
+// CHECK-DAG: @v1
 __device__ int v1;
 
-// NEGCHK-NOT: @v2
+// CHECK-DAG: @v2
 __constant__ int v2;
 
+// Check unused device/constant variables are eliminated.
+
 // NEGCHK-NOT: @_ZL2v3
 constexpr int v3 = 1;
 
 // Check managed variables are always kept.
 
 // CHECK-DAG: @v4
 __managed__ int v4;
 

llvm/lib/Transforms/IPO/Internalize.cpp

bool InternalizePass::shouldPreserveGV(const GlobalValue &GV) {
   // Available externally is really just a "declaration with a body".
   if (GV.hasAvailableExternallyLinkage())
     return true;
 
   // Assume that dllexported symbols are referenced elsewhere
   if (GV.hasDLLExportStorageClass())
     return true;
 
+  // As the name suggests, 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;
 
   // Check some special cases
   if (AlwaysPreserved.count(GV.getName()))
     return true;
 

llvm/test/Transforms/Internalize/externally-initialized.ll

+; 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