[llvm] bb206cb - [NVPTX] Apply global var demotion to private symbols

[Quentin Colombet via llvm-commits]

Author: Quentin Colombet
Date: 2023-08-09T14:41:01+02:00
New Revision: bb206cb131226c375d0e09251e35a3d75dd3f06f

URL: https://github.com/llvm/llvm-project/commit/bb206cb131226c375d0e09251e35a3d75dd3f06f
DIFF: https://github.com/llvm/llvm-project/commit/bb206cb131226c375d0e09251e35a3d75dd3f06f.diff

LOG: [NVPTX] Apply global var demotion to private symbols

When emitting the assembly we perform some late global variables demotion.
Prior to this patch, this optimization was only performed on variables with
the internal linkage whereas any local global variable can be demoted.

Fix that by using `hasLocalLinkage` instead of `hasInternalLinkage`.

Without this change, global variables with the `private` linkage wouldn't
be demoted.

Differential Revision: https://reviews.llvm.org/D154507

Added: 
    llvm/test/CodeGen/NVPTX/demote-vars.ll

Modified: 
    llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 71b70766bf9e13..ccf811f1d71fc8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -674,11 +674,11 @@ static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
  * Currently, this is valid for CUDA shared variables, which have local
  * scope and global lifetime. So the conditions to check are :
  * 1. Is the global variable in shared address space?
- * 2. Does it have internal linkage?
+ * 2. Does it have local linkage?
  * 3. Is the global variable referenced only in one function?
  */
 static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
-  if (!gv->hasInternalLinkage())
+  if (!gv->hasLocalLinkage())
     return false;
   PointerType *Pty = gv->getType();
   if (Pty->getAddressSpace() != ADDRESS_SPACE_SHARED)

diff  --git a/llvm/test/CodeGen/NVPTX/demote-vars.ll b/llvm/test/CodeGen/NVPTX/demote-vars.ll
new file mode 100644
index 00000000000000..80c322a578fc26
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/demote-vars.ll
@@ -0,0 +1,82 @@
+; RUN: llc -o - %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+
+; Check that we do global variable demotion when the symbols don't need to be
+; shared across modules or functions.
+
+; Variable visible globally can't be demoted.
+; CHECK: .visible .shared .align 4 .u64 external_global
+ at external_global = addrspace(3) global i64 undef, align 4
+
+; Not externally visible global used in only one function => demote.
+; It should live in @define_internal_global.
+ at internal_global = internal addrspace(3) global i64 undef, align 4
+
+; Not externally visible global used in several functions => keep.
+; CHECK: .shared .align 4 .u64 internal_global_used_in_
diff erent_fcts
+ at internal_global_used_in_
diff erent_fcts = internal addrspace(3) global i64 undef, align 4
+
+; Not externally visible global used in only one function => demote.
+; It should live in @define_private_global.
+ at private_global = private addrspace(3) global i64 undef, align 4
+
+; Not externally visible global used in only one function
+; (but with several uses) => demote.
+; It should live in @define_private_global_more_than_one_use.
+ at private_global_used_more_than_once_in_same_fct = private addrspace(3) global i64 undef, align 4
+
+; CHECK-LABEL: define_external_global(
+define void @define_external_global(i64 %val) {
+  store i64 %val, ptr addrspace(3) @external_global
+  ret void
+}
+
+; CHECK-LABEL: define_internal_global(
+; Demoted `internal_global` should live here.
+; CHECK: .shared .align 4 .u64 internal_global
+define void @define_internal_global(i64 %val) {
+  store i64 %val, ptr addrspace(3) @internal_global
+  ret void
+}
+
+; CHECK-LABEL: define_internal_global_with_
diff erent_fct1(
+; CHECK-NOT: .shared .align 4 .u64 internal_global_used_in_
diff erent_fcts
+define void @define_internal_global_with_
diff erent_fct1(i64 %val) {
+  store i64 %val, ptr addrspace(3) @internal_global_used_in_
diff erent_fcts
+  ret void
+}
+
+; CHECK-LABEL: define_internal_global_with_
diff erent_fct2(
+; CHECK-NOT: .shared .align 4 .u64 internal_global_used_in_
diff erent_fcts
+define void @define_internal_global_with_
diff erent_fct2(i64 %val) {
+  store i64 %val, ptr addrspace(3) @internal_global_used_in_
diff erent_fcts
+  ret void
+}
+
+; CHECK-LABEL: define_private_global(
+; Demoted `private_global` should live here.
+; CHECK: .shared .align 4 .u64 private_global
+define void @define_private_global(i64 %val) {
+  store i64 %val, ptr addrspace(3) @private_global
+  ret void
+}
+
+; CHECK-LABEL: define_private_global_more_than_one_use(
+; Demoted `private_global_used_more_than_once_in_same_fct` should live here.
+; CHECK: .shared .align 4 .u64 private_global_used_more_than_once_in_same_fct
+;
+; Also check that the if-then is still here, otherwise we may not be testing
+; the "more-than-one-use" part.
+; CHECK: st.shared.u64   [private_global_used_more_than_once_in_same_fct],
+; CHECK: mov.u64 %[[VAR:.*]], 25
+; CHECK: st.shared.u64   [private_global_used_more_than_once_in_same_fct], %[[VAR]]
+define void @define_private_global_more_than_one_use(i64 %val, i1 %cond) {
+  store i64 %val, ptr addrspace(3) @private_global_used_more_than_once_in_same_fct
+  br i1 %cond, label %then, label %end
+
+then:
+  store i64 25, ptr addrspace(3) @private_global_used_more_than_once_in_same_fct
+  br label %end
+
+end:
+  ret void
+}