[llvm] r245785 - [NVPTX] Allow undef value as global initializer

[Jingyue Wu via llvm-commits]

Author: jingyue
Date: Sat Aug 22 00:40:26 2015
New Revision: 245785

URL: http://llvm.org/viewvc/llvm-project?rev=245785&view=rev
Log:
[NVPTX] Allow undef value as global initializer

Summary:
__shared__ variable may now emit undef value as initializer, do not
throw error on that.

Test Plan: test/CodeGen/NVPTX/global-addrspace.ll

Patch by Xuetian Weng

Reviewers: jholewinski, tra, jingyue

Subscribers: llvm-commits, jholewinski

Differential Revision: http://reviews.llvm.org/D12242

Added:
    llvm/trunk/test/CodeGen/NVPTX/global-addrspace.ll
Modified:
    llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=245785&r1=245784&r2=245785&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp Sat Aug 22 00:40:26 2015
@@ -1183,9 +1183,11 @@ void NVPTXAsmPrinter::printModuleLevelGV
           printScalarConstant(Initializer, O);
         }
       } else {
-        // The frontend adds zero-initializer to variables that don't have an
-        // initial value, so skip warning for this case.
-        if (!GVar->getInitializer()->isNullValue()) {
+        // The frontend adds zero-initializer to device and constant variables
+        // that don't have an initial value, and UndefValue to shared
+        // variables, so skip warning for this case.
+        if (!GVar->getInitializer()->isNullValue() &&
+            !isa<UndefValue>(GVar->getInitializer())) {
           report_fatal_error("initial value of '" + GVar->getName() +
                              "' is not allowed in addrspace(" +
                              Twine(PTy->getAddressSpace()) + ")");

Added: llvm/trunk/test/CodeGen/NVPTX/global-addrspace.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/global-addrspace.ll?rev=245785&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/global-addrspace.ll (added)
+++ llvm/trunk/test/CodeGen/NVPTX/global-addrspace.ll Sat Aug 22 00:40:26 2015
@@ -0,0 +1,12 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
+
+; PTX32: .visible .global .align 4 .u32 i;
+; PTX32: .visible .const .align 4 .u32 j;
+; PTX32: .visible .shared .align 4 .u32 k;
+; PTX64: .visible .global .align 4 .u32 i;
+; PTX64: .visible .const .align 4 .u32 j;
+; PTX64: .visible .shared .align 4 .u32 k;
+ at i = addrspace(1) externally_initialized global i32 0, align 4
+ at j = addrspace(4) externally_initialized global i32 0, align 4
+ at k = addrspace(3) global i32 undef, align 4