[PATCH] D12242: [NVPTX] Allow undef value as global initializer

[Xuetian Weng via llvm-commits]

wengxt created this revision.
wengxt added reviewers: jingyue, jholewinski.
wengxt added a subscriber: llvm-commits.
Herald added a subscriber: jholewinski.

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

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

http://reviews.llvm.org/D12242

Files:
  lib/Target/NVPTX/NVPTXAsmPrinter.cpp
  test/CodeGen/NVPTX/global-addrspace.ll

Index: test/CodeGen/NVPTX/global-addrspace.ll
===================================================================
--- /dev/null
+++ test/CodeGen/NVPTX/global-addrspace.ll
@@ -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
Index: lib/Target/NVPTX/NVPTXAsmPrinter.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -1185,7 +1185,8 @@
       } 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()) {
+        if (!GVar->getInitializer()->isNullValue() &&
+            !isa<UndefValue>(GVar->getInitializer())) {
           report_fatal_error("initial value of '" + GVar->getName() +
                              "' is not allowed in addrspace(" +
                              Twine(PTy->getAddressSpace()) + ")");


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D12242.32837.patch
Type: text/x-patch
Size: 1523 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150821/5638bded/attachment.bin>