RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2025-01-31 20:51:52 +01:00
Code Issues Projects Releases Wiki Activity

[NVPTX] Allow undef value as global initializer

Browse Source 
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

llvm-svn: 245785
This commit is contained in:
Jingyue Wu 2015-08-22 05:40:26 +00:00
parent 5f184b4bb7
commit 2549889c79
2 changed files with 17 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
lib/Target/NVPTX/NVPTXAsmPrinter.cpp
View File

@ -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<UndefValue>(GVar->getInitializer())) {
report_fatal_error("initial value of '" + GVar->getName() +
"' is not allowed in addrspace(" +
Twine(PTy->getAddressSpace()) + ")");

12
test/CodeGen/NVPTX/global-addrspace.ll Normal file
View File

@ -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