[GlobalOpt][FIX] Do not embed initializers into AS!=0 globals

[jdoerfert]

Summary

Not all address spaces support initializers for globals and we can
therefore not set them without checking if they are allowed. This
patch adds a hook into TTI to check if an AS allows non-undef
initializers. We disable it for all but address space 0 by default,
NVPTX and AMDGPU targets allow all but address space 3.

[jdoerfert]

clang format

[tra]

For NVPTX, it's the __shared__ AS that's the odd one. Static initializers do work for __device__ and __constant__ globals.
I think this may be something we should plumb via TTI with AS(0) being the only AS allowed by default.

[jdoerfert]

Introduce TTI hook

[tra]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371148)

While we should not have globals with ADDRESS_SPACE_LOCAL or ADDRESS_SPACE_PARAM, it may be worth to return false for them, too.

llvm/test/Transforms/GlobalOpt/address_space_initializer.ll

2

You may want to add another run for NVPTX with non-default AS and check that we do create the initializer.

[jdoerfert]

Better test coverage, added assertion

[jdoerfert]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371148)

That sounds to me more like an assertion. I'll add one for now, you let me know if that is ok too.

[tra]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371148)

Assertions are for compiler programming errors, not for diagnosing invalid user-provided input.
Considering that one can write IR with a global var in an invalid AS I think assertion is not the right choice here.
An error would be more appropriate, IMO.

[jdoerfert]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371148)

TBH, I think such an input is broken and the verifier should trigger, which it does not: https://godbolt.org/z/xr1v361Wo
Assuming the IR is "valid" we should not misuse this interface with an AS that is not allowed for globals, IMHO.
That said, I'm not even sure what "an error" would look like here.

[tra]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371148)

Right. TTI just does not have enough context to tell whether this particular call indicates a problem. Perhaps the best we can do here is to just provide the answer whether the AS can/can't handle global initializers and leave it up to the callers to deal with overall IR correctness.

In this particular patch, returning false for local/param AS would be valid. It will block the optimization and that would the the correct action regardless of soundness of the IR we're optimizing. Improving target-specific IR checks in the verifier would be nice, but it's beyond the scope of this patch and should be done separately.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUTargetTransformInfo.h

187 ↗

(On Diff #371190)

Should also disallow REGION_ADDRESS, and PRIVATE_ADDRESS (we don't even allow private globals)

[jdoerfert]

Return false for nonsensical address spaces

[tra]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371458)

Nit: Making it a switch and enumerating all known AS would make it harder to give a wrong answer if/when a new AS is added.

[jdoerfert]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371458)

The AMD enum has dozens of members :( and we already do these checks in other places. I'd leave it like this for now if it's ok.

[tra]

llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h

52 ↗

(On Diff #371458)

OK. This is fine, too.

[arsenm]

llvm/include/llvm/Analysis/TargetTransformInfoImpl.h

109 ↗

(On Diff #371949)

I'm not sure this should be so strict by default. At the minimum it should probably allow the datalayout's default global address space

[jdoerfert]

llvm/include/llvm/Analysis/TargetTransformInfoImpl.h

109 ↗

(On Diff #371949)

I have no strong feelings either way. This seems more cautious and maybe we want targets to overwrite it explicitly? That said, feel free to patch it.

llvm/lib/Transforms/IPO/GlobalOpt.cpp

if (GS.StoredType == GlobalStatus::StoredOnce && GS.StoredOnceValue) {
     auto *SOVConstant = dyn_cast<Constant>(GS.StoredOnceValue);
     if (SOVConstant && SOVConstant->canTrap())
       return Changed;
 
     // If the initial value for the global was an undef value, and if only
     // one other value was stored into it, we can just change the
     // initializer to be the stored value, then delete all stores to the
     // global.  This allows us to mark it constant.
+    // This is restricted to address space 0 globals. NVPTX, for example,
+    // does not support initializers for shared memory (AS 3).
     if (SOVConstant && SOVConstant->getType() == GV->getValueType() &&
-        isa<UndefValue>(GV->getInitializer())) {
+        isa<UndefValue>(GV->getInitializer()) && GV->getType()->getAddressSpace() == 0) {
       // Change the initial value here.
       GV->setInitializer(SOVConstant);
 
       // Clean up any obviously simplifiable users now.
       CleanupConstantGlobalUsers(GV, GV->getInitializer(), DL, GetTLI);
 
       if (GV->use_empty()) {
         LLVM_DEBUG(dbgs() << "   *** Substituting initializer allowed us to "

llvm/test/Transforms/GlobalOpt/address_space_initializer.ll

+; RUN: opt -passes=globalopt < %s -S | FileCheck %s
+
+; Check that we don't try to set a global initializer for non AS(0) globals.
+
+@g0 = internal global i16 undef
+@g1 = internal addrspace(3) global i16 undef
+; CHECK: internal unnamed_addr constant i16 3
+; CHECK: internal unnamed_addr addrspace(3) global i16 undef
+
+define void @a() {
+  store i16 3, i16* @g0, align 8
+  store i16 5, i16* addrspacecast (i16 addrspace(3)* @g1 to i16*), align 8
+  ret void
+}
+
+define i8 @get0() optnone noinline {
+  %bc = bitcast i16* @g0 to i8*
+  %gep = getelementptr i8, i8* %bc, i64 1
+  %r = load i8, i8* %gep
+  ret i8 %r
+}
+define i8 @get1() {
+  %ac = addrspacecast i16 addrspace(3)* @g1 to i16*
+  %bc = bitcast i16* %ac to i8*
+  %gep = getelementptr i8, i8* %bc, i64 1
+  %r = load i8, i8* %gep
+  ret i8 %r
+}
+