From 8724a428dfd5e78d7865bb01783708e83f9ed128 Mon Sep 17 00:00:00 2001 From: Jingyue Wu Date: Sat, 22 Aug 2015 05:40:26 +0000 Subject: [PATCH] [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 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@245785 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 8 +++++--- test/CodeGen/NVPTX/global-addrspace.ll | 12 ++++++++++++ 2 files changed, 17 insertions(+), 3 deletions(-) create mode 100644 test/CodeGen/NVPTX/global-addrspace.ll diff --git a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp index ffdac61f44d..16afe3cbfb5 100644 --- a/lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ b/lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -1183,9 +1183,11 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar, 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(GVar->getInitializer())) { report_fatal_error("initial value of '" + GVar->getName() + "' is not allowed in addrspace(" + Twine(PTy->getAddressSpace()) + ")"); diff --git a/test/CodeGen/NVPTX/global-addrspace.ll b/test/CodeGen/NVPTX/global-addrspace.ll new file mode 100644 index 00000000000..4da14c7ff4f --- /dev/null +++ b/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; +@i = addrspace(1) externally_initialized global i32 0, align 4 +@j = addrspace(4) externally_initialized global i32 0, align 4 +@k = addrspace(3) global i32 undef, align 4 -- 2.34.1