[NVPTX] Apply global var demotion to private symbols

[qcolombet]

Summary

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.

[guraypp]

Looks great! This work also solves my issue in D154074

[tra]

This may be potentially problematic for CUDA as we may need to refer to the global variable from the host side. E.g. via cudaMemcpyFromSymbol/cudaMemcpyToSymbol.

[guraypp]

In D154507#4486860, @tra wrote:

This may be potentially problematic for CUDA as we may need to refer to the global variable from the host side.

What about if we add shared memory address space check? The shared memory isn't addressable from the host. Something like below:

if (!gv->hasInternalLinkage() || (!gv->hasLocalLinkage() && Pty->getAddressSpace() != ADDRESS_SPACE_SHARED))

[tra]

In D154507#4486877, @guraypp wrote:

In D154507#4486860, @tra wrote:

What about if we add shared memory address space check? The shared memory isn't addressable from the host. Something like below:

if (!gv->hasInternalLinkage() || (!gv->hasLocalLinkage() && Pty->getAddressSpace() != ADDRESS_SPACE_SHARED))

For shared variables demotion should be fine, though I'm not sure what it would buy us. Shared variables, even if they get demoted would still be allocated statically, so effectively it only affects the symbol visibility without saving any resources, which would be the case for non-shared variables.

[tra]

Ah, never mind. We're only dealing with shared variables here. The patch is fine.

[guraypp]

In D154507#4486918, @tra wrote:

Ah, never mind. We're only dealing with shared variables here. The patch is fine.

Awesome thanks!

[mehdi_amini]

In D154507#4486860, @tra wrote:

This may be potentially problematic for CUDA as we may need to refer to the global variable from the host side. E.g. via cudaMemcpyFromSymbol/cudaMemcpyToSymbol.

If something is to be referenced from outside a LLVM module, shouldn't we just avoid making it internal in the first place? Neither internal nor private are supposed to be visible outside the current module I believe.

[tra]

In D154507#4487074, @mehdi_amini wrote:

In D154507#4486860, @tra wrote:

This may be potentially problematic for CUDA as we may need to refer to the global variable from the host side. E.g. via cudaMemcpyFromSymbol/cudaMemcpyToSymbol.

If something is to be referenced from outside a LLVM module, shouldn't we just avoid making it internal in the first place? Neither internal nor private are supposed to be visible outside the current module I believe.

Please ignore, I missed that this demotion only applies to the shared variables which can't be accessed from the host that way.

[jdoerfert]

I don't think this is (and was) correct, it does not account for capturing.

We have a pass that does this in the middle end (https://llvm.org/doxygen/GlobalOpt_8cpp_source.html#l01359), which reminds me of https://discourse.llvm.org/t/rfc-cleaning-up-the-nvidia-and-potentially-amd-gpu-backend/71680.

[mehdi_amini]

I'm a bit confused @jdoerfert : are there two different uses of the word "demoted" here? Seems like you're referring to a transformation that turns a global into an alloca, whereas here is it just about "demoting" the linkage type of the global?

[jdoerfert]

In D154507#4487519, @mehdi_amini wrote:

I'm a bit confused @jdoerfert : are there two different uses of the word "demoted" here? Seems like you're referring to a transformation that turns a global into an alloca, whereas here is it just about "demoting" the linkage type of the global?

I do not see the linkage being changed, if so, to what? Local (=private or internal) is already as "good" as it gets.

Given the test result, and the comment ("Find out if a global variable can be demoted to local scope."), it looks to me as if the scope is changed, thus where the variable is "declared/visible".
That said, I was expecting (w/o knowing for sure), that a variable scoped in a function cannot be accessed by other functions. If that is wrong, the first part of my comment is mood.
I still believe we don't want these special casings in our backends, but that is for later. If the scope doesn't preclude access, this is fine.

EDIT:
I looked at the PTX spec online but could not figure it out right away. One can always try to escape the address, recurse, check that the address is still the same and that outside accesses are visible as they should be.

[mehdi_amini]

Ah I see what you mean! Scratch what I wrote above...

[qcolombet]

Hi @jdoerfert,

In D154507#4487631, @jdoerfert wrote:

In D154507#4487519, @mehdi_amini wrote:

I'm a bit confused @jdoerfert : are there two different uses of the word "demoted" here? Seems like you're referring to a transformation that turns a global into an alloca, whereas here is it just about "demoting" the linkage type of the global?

I do not see the linkage being changed, if so, to what? Local (=private or internal) is already as "good" as it gets.

Given the test result, and the comment ("Find out if a global variable can be demoted to local scope."), it looks to me as if the scope is changed, thus where the variable is "declared/visible".
That said, I was expecting (w/o knowing for sure), that a variable scoped in a function cannot be accessed by other functions. If that is wrong, the first part of my comment is mood.

I think a variable scoped in a function can be accessed by other functions. Meaning, the symbol may not be visible, but the address pointed by this symbol can be accessed (i.e., it is okay to escape), because the allocations are handled at a higher level (the driver).
@guraypp do you have more details on this?

I still believe we don't want these special casings in our backends, but that is for later. If the scope doesn't preclude access, this is fine.

@guraypp could you explain why this optimization is useful?
I don't remember. I would expect the visibility doesn't matter as long as the (meta)data around what needs to be allocated before a kernel launch is accurate, but maybe this info is derived from the symbols visibility?

EDIT:
I looked at the PTX spec online but could not figure it out right away. One can always try to escape the address, recurse, check that the address is still the same and that outside accesses are visible as they should be.

Looking at how usedInOneFunc is implemented, you're right that the address could escape. That being said, I don't think this is an issue since the memory allocation is handled at the kernel level. Anyhow, this is orthogonal to this patch IMHO.

Cheers,
-Quentin

[qcolombet]

In D154507#4509642, @qcolombet wrote:

Hi @jdoerfert,

In D154507#4487631, @jdoerfert wrote:

In D154507#4487519, @mehdi_amini wrote:

I'm a bit confused @jdoerfert : are there two different uses of the word "demoted" here? Seems like you're referring to a transformation that turns a global into an alloca, whereas here is it just about "demoting" the linkage type of the global?

I do not see the linkage being changed, if so, to what? Local (=private or internal) is already as "good" as it gets.

Given the test result, and the comment ("Find out if a global variable can be demoted to local scope."), it looks to me as if the scope is changed, thus where the variable is "declared/visible".
That said, I was expecting (w/o knowing for sure), that a variable scoped in a function cannot be accessed by other functions. If that is wrong, the first part of my comment is mood.

I think a variable scoped in a function can be accessed by other functions. Meaning, the symbol may not be visible, but the address pointed by this symbol can be accessed (i.e., it is okay to escape), because the allocations are handled at a higher level (the driver).
@guraypp do you have more details on this?

Spoke with @guraypp offline and he confirms that was the case.

I still believe we don't want these special casings in our backends, but that is for later. If the scope doesn't preclude access, this is fine.

@guraypp could you explain why this optimization is useful?
I don't remember. I would expect the visibility doesn't matter as long as the (meta)data around what needs to be allocated before a kernel launch is accurate, but maybe this info is derived from the symbols visibility?

@guraypp checked on Hopper and the driver seems to do the right allocation with or without the demotion. (Meaning the demotion doesn't seem that useful anymore.)
That being said, it may still be useful for older driver.
Anyhow, scope doesn't preclude access so I think @jdoerfert, you're fine with the patch, right?

EDIT:
I looked at the PTX spec online but could not figure it out right away. One can always try to escape the address, recurse, check that the address is still the same and that outside accesses are visible as they should be.

Looking at how usedInOneFunc is implemented, you're right that the address could escape. That being said, I don't think this is an issue since the memory allocation is handled at the kernel level. Anyhow, this is orthogonal to this patch IMHO.

Cheers,
-Quentin

[qcolombet]

Ping @jdoerfert

Does the patch work for you with the additional context?
The TL;DR is the semantic is correct but more recent drivers may not need that optimization anymore.

[jdoerfert]

In D154507#4531775, @qcolombet wrote:

Ping @jdoerfert

Does the patch work for you with the additional context?
The TL;DR is the semantic is correct but more recent drivers may not need that optimization anymore.

SG, thx.

llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
 
   return true;
 }
 
 /* Find out if a global variable can be demoted to local scope.
  * 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)
     return false;
 
   const Function *oneFunc = nullptr;
 
   bool flag = usedInOneFunc(gv, oneFunc);

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

+; 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
+@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.
+@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_different_fcts
+@internal_global_used_in_different_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.
+@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.
+@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_different_fct1(
+; CHECK-NOT: .shared .align 4 .u64 internal_global_used_in_different_fcts
+define void @define_internal_global_with_different_fct1(i64 %val) {
+  store i64 %val, ptr addrspace(3) @internal_global_used_in_different_fcts
+  ret void
+}
+
+; CHECK-LABEL: define_internal_global_with_different_fct2(
+; CHECK-NOT: .shared .align 4 .u64 internal_global_used_in_different_fcts
+define void @define_internal_global_with_different_fct2(i64 %val) {
+  store i64 %val, ptr addrspace(3) @internal_global_used_in_different_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
+}