[AMDGPU] Add GlobalDCE before internalization pass

[yaxunl]

Summary

The internalization pass only internalizes global variables
with no users. If the global variable has some dead user,
the internalization pass will not internalize it.

To be able to internalize global variables with dead
users, a global dce pass is needed before the
internalization pass.

This patch adds that.

[tra]

Shouldn't the unused ASCs be eliminated by DCE? We seem to be cleaning up the consequences of a problem that happened somewhere else.

[tra]

I do not see any hanging ASCs in the generated IR with clang@HEAD: https://godbolt.org/z/TE4Yhr

What do I miss?

[yaxunl]

In D98783#2632153, @tra wrote:

I do not see any hanging ASCs in the generated IR with clang@HEAD: https://godbolt.org/z/TE4Yhr

What do I miss?

The unused casts are invisible since they are not used by any visible IR's. To see them, you need to dump the users of these global variables explicitly.

[yaxunl]

In D98783#2632143, @tra wrote:

Shouldn't the unused ASCs be eliminated by DCE? We seem to be cleaning up the consequences of a problem that happened somewhere else.

DCE does not eliminate these unused ASCs because normally they should not exist. That's why we do not want them to be kept in the IR emitted by clang.

[yaxunl]

bug fix. only remove unused casts.

[tra]

In D98783#2632244, @yaxunl wrote:

In D98783#2632143, @tra wrote:

Shouldn't the unused ASCs be eliminated by DCE? We seem to be cleaning up the consequences of a problem that happened somewhere else.

DCE does not eliminate these unused ASCs because normally they should not exist. That's why we do not want them to be kept in the IR emitted by clang.

TBH, these 'invisible' ASCs do bother me. I wonder what else is present in the IR that we can't easily examine with existing tools?
In principle I'm OK with eliminating such dangling ASCs here as a short-term workaround, but I think it's potentially a more general issue which needs a more general solution.

One way to consistently deal with something like that is to codegen something using those ASCs, but which would be considered unused and later DCE'd along with unused ASCs.
We could use some sort of counterpart for @llvm.used only for potentially unused IR we create.
Tying them to such @llvm_maybe_unused would avoid the problem.

[yaxunl]

In D98783#2633261, @tra wrote:

In D98783#2632244, @yaxunl wrote:

In D98783#2632143, @tra wrote:

Shouldn't the unused ASCs be eliminated by DCE? We seem to be cleaning up the consequences of a problem that happened somewhere else.

DCE does not eliminate these unused ASCs because normally they should not exist. That's why we do not want them to be kept in the IR emitted by clang.

TBH, these 'invisible' ASCs do bother me. I wonder what else is present in the IR that we can't easily examine with existing tools?
In principle I'm OK with eliminating such dangling ASCs here as a short-term workaround, but I think it's potentially a more general issue which needs a more general solution.

The invisible LLVM constants only happen when they are created but not used later, which is rare since usually a constant is created and used immediately, making them visible.

One way to consistently deal with something like that is to codegen something using those ASCs, but which would be considered unused and later DCE'd along with unused ASCs.
We could use some sort of counterpart for @llvm.used only for potentially unused IR we create.
Tying them to such @llvm_maybe_unused would avoid the problem.

That's one solution. Another solution is to add a member function to LLVM module to clean up unused LLVM constants and add check to LLVM IR verifier to make sure there is no unused constants.

[tra]

We may want to add someone with more expertise with the IR as a reviewer. I'd like an educated opinion on whether the invisible dangling IR is something that needs fixing in general or if it's OK to just clean it up in this particular case. Or both.

@rjmccall, @rsmith -- do you have any suggestions -- either on the subject of the invisible dangling IR or on who may be the right person to talk to?

[rjmccall]

In D98783#2642269, @tra wrote:

We may want to add someone with more expertise with the IR as a reviewer. I'd like an educated opinion on whether the invisible dangling IR is something that needs fixing in general or if it's OK to just clean it up in this particular case. Or both.

@rjmccall, @rsmith -- do you have any suggestions -- either on the subject of the invisible dangling IR or on who may be the right person to talk to?

If Clang is creating constants unnecessarily, we should try to avoid that on general compile time / memory usage grounds unless doing so is a serious problem. But I don't think it should be our responsibility to GC unused constant expressions in order to make DCE work.

It's extremely common to create bitcast global constants. So the fact that this happens with addrspacecasts but hasn't been a persistent problem with bitcasts makes more suspect that DCE actually tries to handle this, but something about what it's doing isn't aware of addrspacecast. And indeed, LLVM's GlobalDCE seems to call a method called removeDeadConstantUsers() before concluding that a constant can't be thrown away. So either you're using a different transform that needs to do the same thing, or something is stopping removeDeadConstantUsers() from eliminating this addrspacecast.

[yaxunl]

In D98783#2657656, @rjmccall wrote:

In D98783#2642269, @tra wrote:

We may want to add someone with more expertise with the IR as a reviewer. I'd like an educated opinion on whether the invisible dangling IR is something that needs fixing in general or if it's OK to just clean it up in this particular case. Or both.

@rjmccall, @rsmith -- do you have any suggestions -- either on the subject of the invisible dangling IR or on who may be the right person to talk to?

If Clang is creating constants unnecessarily, we should try to avoid that on general compile time / memory usage grounds unless doing so is a serious problem. But I don't think it should be our responsibility to GC unused constant expressions in order to make DCE work.

It's extremely common to create bitcast global constants. So the fact that this happens with addrspacecasts but hasn't been a persistent problem with bitcasts makes more suspect that DCE actually tries to handle this, but something about what it's doing isn't aware of addrspacecast. And indeed, LLVM's GlobalDCE seems to call a method called removeDeadConstantUsers() before concluding that a constant can't be thrown away. So either you're using a different transform that needs to do the same thing, or something is stopping removeDeadConstantUsers() from eliminating this addrspacecast.

You are right. GlobalDCE is able to remove the dead addr space casts added by clang. The reason why the useless global variables are not internalized is that the internalization pass is done before GlobalDCE, whereas the internalization pass only internalizes useless global variables. To be able to internalize global variables with dead users, I need to add another GlobalDCE before internalization pass.

[yaxunl]

Add GlobalDCE before internalization pass.

[tra]

LGTM with a test nit.

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

24–30

Mixing CHECK and CHECK-NOT is tricky and, in general, only works if things are always in the same order.

E.g. if does v3 get emitted after v4, the test will still pass.

One way to deal with that would be to split the positive and negative checks into separate runs.
First one would check the variables we do want to keep with CHECK-DAG.
The other one would only check for the absence of the variables with -NOT.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

581

Should we move where the internalize pass is added instead?

[yaxunl]

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

24–30

will do

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

581

we need two global dce, one before internalize, one after internalize.

the first global dce will eliminate the dead users of the global vars, but not the vars themselves since they are external.

the internalize pass will make the useless vars internal.

the second global dec will eliminate the useless internal global vars. Without the second global dce, the useless global var will not be eliminated.

[yaxunl]

revised tests by Artem's comments.

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

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
+// RUN:   -target-cpu gfx906 | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
+// RUN:   -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
+// RUN:   -target-cpu gfx906 | FileCheck -check-prefix=NEGCHK %s
+
+#include "Inputs/cuda.h"
+
+// 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
+__device__ int v1;
+
+// NEGCHK-NOT: @v2
+__constant__ int v2;
+
+// NEGCHK-NOT: @_ZL2v3
+constexpr int v3 = 1;
+
+// Check managed variables are always kept.
+
+// CHECK-DAG: @v4
+__managed__ int v4;
+
+// Check used device/constant variables are not eliminated.
+// CHECK-DAG: @u1
+__device__ int u1;
+
+// CHECK-DAG: @u2
+__constant__ int u2;
+
+// Check u3 is kept because its address is taken.
+// CHECK-DAG: @_ZL2u3
+constexpr int u3 = 2;
+
+// Check u4 is not kept because it is not ODR-use.
+// NEGCHK-NOT: @_ZL2u4
+constexpr int u4 = 3;
+
+__device__ int fun1(const int& x);
+
+__global__ void kern1(int *x) {
+  *x = u1 + u2 + fun1(u3) + u4;
+}

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

PB.registerPipelineEarlySimplificationEPCallback(
       [this](ModulePassManager &PM, PassBuilder::OptimizationLevel Level) {
         if (Level == PassBuilder::OptimizationLevel::O0)
           return;
 
         PM.addPass(AMDGPUUnifyMetadataPass());
         PM.addPass(AMDGPUPrintfRuntimeBindingPass());
 
         if (InternalizeSymbols) {
+          // Global variables may have dead uses which need to be removed.
+          // Otherwise these useless global variables will not get internalized.
+          PM.addPass(GlobalDCEPass());
           PM.addPass(InternalizePass(mustPreserveGV));
         }
         PM.addPass(AMDGPUPropagateAttributesLatePass(*this));
         if (InternalizeSymbols) {
           PM.addPass(GlobalDCEPass());
         }
         if (EarlyInlineAll && !EnableFunctionCalls)
           PM.addPass(AMDGPUAlwaysInlinePass());
       });