summaryrefslogtreecommitdiffstats
path: root/clang/lib/CodeGen/CodeGenModule.cpp
Commit message (Collapse)AuthorAgeFilesLines
* PR25910: clang allows two var definitions with the same mangled nameAndrey Bokhanko2016-01-141-27/+86
| | | | | | | | | | | Proper diagnostic and resolution of mangled names' conflicts in variables. When there is a declaration and a definition using the same name but different types, we emit what is in the definition. When there are two conflicting definitions, we issue an error. Differential Revision: http://reviews.llvm.org/D15686 llvm-svn: 257754
* [MS ABI] Complete and base constructor GlobalDecls must have the same nameDavid Majnemer2016-01-081-1/+14
| | | | | | | | | | | | | | | | | Clang got itself into the situation where we mangled the same constructor twice with two different constructor types. After one of the constructors were utilized, the tag used for one of the types changed from class to struct because a class template became complete. This resulted in one of the constructor types varying from the other constructor. Instead, force "base" constructor types to "complete" if the ABI doesn't have constructor variants. This will ensure that GlobalDecls for both variants will get the same mangled name. This fixes PR26029. llvm-svn: 257205
* [Driver] Add support for -fno-builtin-foo options.Chad Rosier2016-01-061-1/+2
| | | | | | | Addresses PR4941 and rdar://6756912. http://reviews.llvm.org/D15195 llvm-svn: 256937
* [OpenMP] Reapply rL256842: [OpenMP] Offloading descriptor registration and ↵Samuel Antao2016-01-061-0/+12
| | | | | | | | | | | | device codegen. This patch attempts to fix the regressions identified when the patch was committed initially. Thanks to Michael Liao for identifying the fix in the offloading metadata generation related with side effects in evaluation of function arguments. llvm-svn: 256933
* [OpenMP] Revert rL256842: [OpenMP] Offloading descriptor registration and ↵Samuel Antao2016-01-051-12/+0
| | | | | | | | device codegen. It was causing two regression, so I'm reverting until the cause is found. llvm-svn: 256858
* [OpenMP] Offloading descriptor registration and device codegen.Samuel Antao2016-01-051-0/+12
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: In order to offloading work properly two things need to be in place: - a descriptor with all the offloading information (device entry functions, and global variable) has to be created by the host and registered in the OpenMP offloading runtime library. - all the device functions need to be emitted for the device and a convention has to be in place so that the runtime library can easily map the host ID of an entry point with the actual function in the device. This patch adds support for these two things. However, only entry functions are being registered given that 'declare target' directive is not yet implemented. About offloading descriptor: The details of the descriptor are explained with more detail in http://goo.gl/L1rnKJ. Basically the descriptor will have fields that specify the number of devices, the pointers to where the device images begin and end (that will be defined by the linker), and also pointers to a the begin and end of table whose entries contain information about a specific entry point. Each entry has the type: ``` struct __tgt_offload_entry{ void *addr; char *name; int64_t size; }; ``` and will be implemented in a pre determined (ELF) section `.omp_offloading.entries` with 1-byte alignment, so that when all the objects are linked, the table is in that section with no padding in between entries (will be like a C array). The code generation ensures that all `__tgt_offload_entry` entries are emitted in the same order for both host and device so that the runtime can have the corresponding entries in both host and device in same index of the table, and efficiently implement the mapping. The resulting descriptor is registered/unregistered with the runtime library using the calls `__tgt_register_lib` and `__tgt_unregister_lib`. The registration is implemented in a high priority global initializer so that the registration happens always before any initializer (that can potentially include target regions) is run. The driver flag -omptargets= was created to specify a comma separated list of devices the user wants to support so that the new functionality can be exercised. Each device is specified with its triple. About target codegen: The target codegen is pretty much straightforward as it reuses completely the logic of the host version for the same target region. The tricky part is to identify the meaningful target regions in the device side. Unlike other programming models, like CUDA, there are no already outlined functions with attributes that mark what should be emitted or not. So, the information on what to emit is passed in the form of metadata in host bc file. This requires a new option to pass the host bc to the device frontend. Then everything is similar to what happens in CUDA: the global declarations emission is intercepted to check to see if it is an "interesting" declaration. The difference is that instead of checking an attribute, the metadata information in checked. Right now, there is only a form of metadata to pass information about the device entry points (target regions). A class `OffloadEntriesInfoManagerTy` was created to manage all the information and queries related with the metadata. The metadata looks like this: ``` !omp_offload.info = !{!0, !1, !2, !3, !4, !5, !6} !0 = !{i32 0, i32 52, i32 77426347, !"_ZN2S12r1Ei", i32 479, i32 13, i32 4} !1 = !{i32 0, i32 52, i32 77426347, !"_ZL7fstatici", i32 461, i32 11, i32 5} !2 = !{i32 0, i32 52, i32 77426347, !"_Z9ftemplateIiET_i", i32 444, i32 11, i32 6} !3 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 99, i32 11, i32 0} !4 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 272, i32 11, i32 3} !5 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 127, i32 11, i32 1} !6 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 159, i32 11, i32 2} ``` The fields in each metadata entry are (in sequence): Entry 1) an ID of the type of metadata - right now only zero is used meaning "OpenMP target region". Entry 2) a unique ID of the device where the input source file that contain the target region lives. Entry 3) a unique ID of the file where the input source file that contain the target region lives. Entry 4) a mangled name of the function that encloses the target region. Entries 5) and 6) line and column number where the target region was found. Entry 7) is the order the entry was emitted. Entry 2) and 3) are required to distinguish files that have the same function name. Entry 4) is required to distinguish different instances of the same declaration (usually templated ones) Entries 5) and 6) are required to distinguish the particular target region in body of the function (it is possible that a given target region is not an entry point - if clause can evaluate always to zero - and therefore we need to identify the "interesting" target regions. ) This patch replaces http://reviews.llvm.org/D12306. Reviewers: ABataev, hfinkel, tra, rjmccall, sfantao Subscribers: FBrygidyn, piotr.rak, Hahnfeld, cfe-commits Differential Revision: http://reviews.llvm.org/D12614 llvm-svn: 256842
* Attach maximum function count to Module when using PGO mode.Easwaran Raman2015-12-171-2/+5
| | | | | | | | This sets the maximum entry count among all functions in the program to the module using module flags. This allows the optimizer to use this information. Differential Revision: http://reviews.llvm.org/D15163 llvm-svn: 255918
* Cross-DSO control flow integrity (Clang part).Evgeniy Stepanov2015-12-151-18/+84
| | | | | | | | | | | | | | Clang-side cross-DSO CFI. * Adds a command line flag -f[no-]sanitize-cfi-cross-dso. * Links a runtime library when enabled. * Emits __cfi_slowpath calls is bitset test fails. * Emits extra hash-based bitsets for external CFI checks. * Sets a module flag to enable __cfi_check generation during LTO. This mode does not yet support diagnostics. llvm-svn: 255694
* [WinEH] Update clang to use operand bundles on call sitesDavid Majnemer2015-12-151-2/+7
| | | | | | | | | | | This updates clang to use bundle operands to associate an invoke with the funclet which it is contained within. Depends on D15517. Differential Revision: http://reviews.llvm.org/D15518 llvm-svn: 255675
* Revert r254647.Easwaran Raman2015-12-121-3/+0
| | | | | | | | Reason: The testcase fails in many architectures. Differential Revision: http://reviews.llvm.org/D15163 llvm-svn: 255416
* Attach maximum function count to Module when using PGO modeEaswaran Raman2015-12-121-0/+3
| | | | | | | | | This sets the maximum entry count among all functions in the program to the module using module flags. This allows the optimizer to use this information. Differential Revision: http://reviews.llvm.org/D15163 llvm-svn: 255397
* Revert "[x86] Exclusion of incorrect include headers paths for MCU target"Reid Kleckner2015-12-051-52/+17
| | | | | | | | | | This reverts commit r254195. From the description, I suspect that the wrong patch was committed here, and this is causing assertion failures in EmitDeferred() when the global value ends up being a bitcast of a global. llvm-svn: 254823
* Add the `pass_object_size` attribute to clang.George Burgess IV2015-12-021-2/+5
| | | | | | | | | | | | | `pass_object_size` is our way of enabling `__builtin_object_size` to produce high quality results without requiring inlining to happen everywhere. A link to the design doc for this attribute is available at the Differential review link below. Differential Revision: http://reviews.llvm.org/D13263 llvm-svn: 254554
* Fix use-after-free when a C++ thread_local variable gets replaced (because itsRichard Smith2015-12-011-2/+2
| | | | | | | type changes when the initializer is attached). Don't hold onto the GlobalVariable*; recompute it from the VarDecl* instead. llvm-svn: 254359
* [x86] Exclusion of incorrect include headers paths for MCU targetAndrey Bokhanko2015-11-271-17/+52
| | | | | | | | Exclusion of /usr/include and /usr/local/include headers paths for MCU target. Differential Revision: http://reviews.llvm.org/D14954 llvm-svn: 254195
* [TLS on Darwin] treat all Darwin platforms in the same way.Manman Ren2015-11-111-1/+1
| | | | | | rdar://problem/9001553 llvm-svn: 252820
* Extract out a function onto CodeGenModule for getting the map ofEric Christopher2015-11-111-0/+29
| | | | | | | features for a particular function, then use it to clean up some code. llvm-svn: 252819
* [TLS on Darwin] change how we handle globals with linkonce or weak linkage.Manman Ren2015-11-111-4/+9
| | | | | | | | | | | | This is about how we handle static member of a template. Before this commit, we use internal linkage for the IR thread-local variable, which is inefficient. With this commit, we will start to follow Itanium C++ ABI. rdar://problem/23415206 Reviewed by John McCall. llvm-svn: 252814
* CodeGen: Remove implicit ilist iterator conversions, NFCDuncan P. N. Exon Smith2015-11-061-1/+2
| | | | | | | Make ilist iterator conversions explicit in clangCodeGen. Eventually I'll remove them everywhere. llvm-svn: 252358
* Fix crash in EmitDeclMetadata modeKeno Fischer2015-11-051-2/+4
| | | | | | | | | | | | | | | | | Summary: This fixes a bug that's easily encountered in LLDB (https://llvm.org/bugs/show_bug.cgi?id=22875). The problem here is that we mangle a name during debug info emission, but never actually emit the actual Decl, so we run into problems in EmitDeclMetadata (which assumes such a Decl exists). Fix that by just skipping metadata emissions for mangled names that don't have associated Decls. Reviewers: rjmccall Subscribers: labath, cfe-commits Differential Revision: http://reviews.llvm.org/D13959 llvm-svn: 252229
* Watch and TV OS: wire up basic ABI choicesTim Northover2015-10-301-0/+2
| | | | | | | This sets the mostly expected Darwin default ABI options for these two platforms. Active changes from these defaults for watchOS are in a later patch. llvm-svn: 251708
* Unify the ObjC entrypoint caches.John McCall2015-10-211-7/+5
| | | | llvm-svn: 250918
* [CodeGen] Remove dead code. NFC.Benjamin Kramer2015-10-151-6/+0
| | | | llvm-svn: 250418
* [CodeGen] [CodeGen] Attach function attributes to functions created inAkira Hatanaka2015-10-081-4/+5
| | | | | | | | | | | | | | | | | | CGBlocks.cpp. This commit fixes a bug in clang's code-gen where it creates the following functions but doesn't attach function attributes to them: __copy_helper_block_ __destroy_helper_block_ __Block_byref_object_copy_ __Block_byref_object_dispose_ rdar://problem/20828324 Differential Revision: http://reviews.llvm.org/D13525 llvm-svn: 249735
* [CodeGen] Check if the Decl pointer passed is null, and if so, returnAkira Hatanaka2015-10-081-7/+15
| | | | | | | | | | | early. This is needed in a patch I plan to commit later, in which a null Decl pointer is passed to SetLLVMFunctionAttributesForDefinition. Relevant discussion is in http://reviews.llvm.org/D13525. llvm-svn: 249722
* [MSVC Compat] Enable ABI impacting non-conforming behavior independently of ↵David Majnemer2015-10-081-1/+1
| | | | | | | | | | -fms-compatibility No ABI for C++ currently makes it possible to implement the standard 100% perfectly. We wrongly hid some of our compatible behavior behind -fms-compatibility instead of tying it to the compiler ABI. llvm-svn: 249656
* Replace double-negated !SourceLocation.isInvalid() with ↵Yaron Keren2015-10-031-1/+1
| | | | | | SourceLocation.isValid(). llvm-svn: 249228
* Use llvm::makeArrayRef. NFC.Craig Topper2015-09-271-1/+1
| | | | llvm-svn: 248678
* Remove attributes minsize and optsize, which conflict with optnone.Akira Hatanaka2015-09-211-4/+2
| | | | | | | | | | | This commit fixes an assert that is triggered when optnone is being added to an IR function that is already marked with minsize and optsize. rdar://problem/22723716 Differential Revision: http://reviews.llvm.org/D13004 llvm-svn: 248191
* Using MD_invariant_groupPiotr Padlewski2015-09-171-1/+1
| | | | | | http://reviews.llvm.org/D12927 llvm-svn: 247933
* [WinEH] Pass the catch adjectives to catchpad directlyReid Kleckner2015-09-161-6/+0
| | | | | | | | | This avoids building a fake LLVM IR global variable just to ferry an i32 down into LLVM codegen. It also puts a nail in the coffin of using MS ABI C++ EH with landingpads, since now we'll assert in the lpad code when flags are present. llvm-svn: 247843
* Decorating vptr load & stores with !invariant.groupPiotr Padlewski2015-09-151-3/+13
| | | | | | | | | | Adding !invariant.group to vptr load/stores for devirtualization purposes. For more goto: http://lists.llvm.org/pipermail/cfe-dev/2015-July/044227.html http://reviews.llvm.org/D12026 llvm-svn: 247725
* Added llvm.module flag for strict vtable pointersPiotr Padlewski2015-09-151-0/+16
| | | | | | | | | | | | It is dangerous to do LTO on code with strict-vtable-pointers, because one module has invariant.group.barriers, and the other one not. In the future I want to just strip all invariant.group metadata from vptrs loads/stores and get rid of invariant.group.barrier calls. http://reviews.llvm.org/D12580 llvm-svn: 247724
* Revert "Always_inline codegen rewrite" and 2 follow-ups.Evgeniy Stepanov2015-09-141-108/+1
| | | | | | | | | | Revert "Update cxx-irgen.cpp test to allow signext in alwaysinline functions." Revert "[CodeGen] Remove wrapper-free always_inline functions from COMDATs" Revert "Always_inline codegen rewrite." Reason for revert: PR24793. llvm-svn: 247620
* [opaque pointer type] Fix a few uses of PointerType::getElementType in favor ↵David Blaikie2015-09-141-2/+1
| | | | | | | | | | of uses of types already available elsewhere These are a few cleanups I happened to have from trying to go in a different direction recently, so just flushing them out while I have them. llvm-svn: 247593
* Revert "[opaque pointer type] update for LLVM API change"David Blaikie2015-09-141-1/+1
| | | | | | | | | | This was the wrong direction to take anyway (because ultimately the GlobalValue needed the pointee type again and /it/ used PointerType::getElementType eventually anyway)... let's go a different way. This reverts commit r236161. llvm-svn: 247586
* [CodeGen] Remove wrapper-free always_inline functions from COMDATsDavid Majnemer2015-09-121-1/+5
| | | | | | always_inline functions without a wrapper don't need to be in a COMDAT. llvm-svn: 247500
* Always_inline codegen rewrite.Evgeniy Stepanov2015-09-121-1/+104
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Current implementation may end up emitting an undefined reference for an "inline __attribute__((always_inline))" function by generating an "available_externally alwaysinline" IR function for it and then failing to inline all the calls. This happens when a call to such function is in dead code. As the inliner is an SCC pass, it does not process dead code. Libc++ relies on the compiler never emitting such undefined reference. With this patch, we emit a pair of 1. internal alwaysinline definition (called F.alwaysinline) 2a. A stub F() { musttail call F.alwaysinline } -- or, depending on the linkage -- 2b. A declaration of F. The frontend ensures that F.inlinefunction is only used for direct calls, and the stub is used for everything else (taking the address of the function, really). Declaration (2b) is emitted in the case when "inline" is meant for inlining only (like __gnu_inline__ and some other cases). This approach, among other nice properties, ensures that alwaysinline functions are always internal, making it impossible for a direct call to such function to produce an undefined symbol reference. This patch is based on ideas by Chandler Carruth and Richard Smith. llvm-svn: 247494
* Revert "Specify target triple in alwaysinline tests."Evgeniy Stepanov2015-09-111-98/+1
| | | | | | | | | Revert "Always_inline codegen rewrite." Breaks gdb & lldb tests. Breaks on Fedora 22 x86_64. llvm-svn: 247491
* Always_inline codegen rewrite.Evgeniy Stepanov2015-09-111-1/+98
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Current implementation may end up emitting an undefined reference for an "inline __attribute__((always_inline))" function by generating an "available_externally alwaysinline" IR function for it and then failing to inline all the calls. This happens when a call to such function is in dead code. As the inliner is an SCC pass, it does not process dead code. Libc++ relies on the compiler never emitting such undefined reference. With this patch, we emit a pair of 1. internal alwaysinline definition (called F.alwaysinline) 2a. A stub F() { musttail call F.alwaysinline } -- or, depending on the linkage -- 2b. A declaration of F. The frontend ensures that F.inlinefunction is only used for direct calls, and the stub is used for everything else (taking the address of the function, really). Declaration (2b) is emitted in the case when "inline" is meant for inlining only (like __gnu_inline__ and some other cases). This approach, among other nice properties, ensures that alwaysinline functions are always internal, making it impossible for a direct call to such function to produce an undefined symbol reference. This patch is based on ideas by Chandler Carruth and Richard Smith. llvm-svn: 247465
* [CUDA] Allow trivial constructors as initializer for __shared__ variables.Artem Belevich2015-09-101-2/+4
| | | | | | Differential Revision: http://reviews.llvm.org/D12739 llvm-svn: 247307
* CFI: Introduce -fsanitize=cfi-icall flag.Peter Collingbourne2015-09-101-5/+15
| | | | | | | | | | This flag causes the compiler to emit bit set entries for functions as well as runtime bitset checks at indirect call sites. Depends on the new function bitset mechanism. Differential Revision: http://reviews.llvm.org/D11857 llvm-svn: 247238
* CodeGen: Introduce CodeGenModule::CreateMetadataIdentifierForType.Peter Collingbourne2015-09-081-0/+19
| | | | | | | | | This function can be used to create a metadata identifier for a specific type. No functionality change, but this will be used by D11857 and D12026. Differential Revision: http://reviews.llvm.org/D12038 llvm-svn: 247098
* Compute and preserve alignment more faithfully in IR-generation.John McCall2015-09-081-41/+53
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Introduce an Address type to bundle a pointer value with an alignment. Introduce APIs on CGBuilderTy to work with Address values. Change core APIs on CGF/CGM to traffic in Address where appropriate. Require alignments to be non-zero. Update a ton of code to compute and propagate alignment information. As part of this, I've promoted CGBuiltin's EmitPointerWithAlignment helper function to CGF and made use of it in a number of places in the expression emitter. The end result is that we should now be significantly more correct when performing operations on objects that are locally known to be under-aligned. Since alignment is not reliably tracked in the type system, there are inherent limits to this, but at least we are no longer confused by standard operations like derived-to-base conversions and array-to-pointer decay. I've also fixed a large number of bugs where we were applying the complete-object alignment to a pointer instead of the non-virtual alignment, although most of these were hidden by the very conservative approach we took with member alignment. Also, because IRGen now reliably asserts on zero alignments, we should no longer be subject to an absurd but frustrating recurring bug where an incomplete type would report a zero alignment and then we'd naively do a alignmentAtOffset on it and emit code using an alignment equal to the largest power-of-two factor of the offset. We should also now be emitting much more aggressive alignment attributes in the presence of over-alignment. In particular, field access now uses alignmentAtOffset instead of min. Several times in this patch, I had to change the existing code-generation pattern in order to more effectively use the Address APIs. For the most part, this seems to be a strict improvement, like doing pointer arithmetic with GEPs instead of ptrtoint. That said, I've tried very hard to not change semantics, but it is likely that I've failed in a few places, for which I apologize. ABIArgInfo now always carries the assumed alignment of indirect and indirect byval arguments. In order to cut down on what was already a dauntingly large patch, I changed the code to never set align attributes in the IR on non-byval indirect arguments. That is, we still generate code which assumes that indirect arguments have the given alignment, but we don't express this information to the backend except where it's semantically required (i.e. on byvals). This is likely a minor regression for those targets that did provide this information, but it'll be trivial to add it back in a later patch. I partially punted on applying this work to CGBuiltin. Please do not add more uses of the CreateDefaultAligned{Load,Store} APIs; they will be going away eventually. llvm-svn: 246985
* Don't crash on a self-alias declarationHal Finkel2015-09-041-0/+5
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | We were crashing in CodeGen given input like this: int self_alias(void) __attribute__((weak, alias("self_alias"))); such a self-alias is invalid, but instead of diagnosing the situation, we'd proceed to produce IR for both the function declaration and the alias. Because we already had a function named 'self_alias', the alias could not be named the same thing, and so LLVM would pick a different name ('self_alias1' for example) for that value. When we later called CodeGenModule::checkAliases, we'd look up the IR value corresponding to the alias name, find the function declaration instead, and then assert in a cast to llvm::GlobalAlias. The easiest way to prevent this is simply to avoid creating the wrongly-named alias value in the first place and issue the diagnostic there (instead of in checkAliases). We detect a related cycle case in CodeGenModule::EmitAliasDefinition already, so this just adds a second such check. Even though the other test cases for this 'alias definition is part of a cycle' diagnostic are in test/Sema/attr-alias-elf.c, I've added a separate regression test for this case. This is because I can't add this check to test/Sema/attr-alias-elf.c without disturbing the other test cases in that file. In order to avoid construction of the bad IR values, this diagnostic is emitted from within CodeGenModule::EmitAliasDefinition (and the relevant declaration is not added to the Aliases vector). The other cycle checks are done within the CodeGenModule::checkAliases function based on the Aliases vector, called from CodeGenModule::Release. However, if there have been errors earlier, HandleTranslationUnit does not call Release, and so checkAliases is never called, and so none of the other diagnostics would be produced. Fixes PR23509. llvm-svn: 246882
* [WebAssembly] Initial WebAssembly support in clangDan Gohman2015-09-031-3/+9
| | | | | | | | | | This implements basic support for compiling (though not yet assembling or linking) for a WebAssembly target. Note that ABI details are not yet finalized, and may change. Differential Revision: http://reviews.llvm.org/D12002 llvm-svn: 246814
* PR17829: Proper diagnostic of mangled names conflictsAndrey Bokhanko2015-08-311-71/+123
| | | | | | | | | | Proper diagnostic and resolution of mangled names conflicts between C++ methods and C functions. This patch implements support for functions/methods only; support for variables is coming separately. Differential Revision: http://reviews.llvm.org/D11297 llvm-svn: 246438
* Allow TLS vars in dllimport/export functions; only inline dllimport ↵Hans Wennborg2015-08-281-0/+38
| | | | | | | | | | | | | | | | | | | | | functions when safe (PR24593) This patch does two things: 1) Don't error about dllimport/export on thread-local static local variables. We put those attributes on static locals in dllimport/export functions implicitly in case the function gets inlined. Now, for TLS variables this is a problem because we can't import such variables, but it's a benign problem becase: 2) Make sure we never inline a dllimport function TLS static locals. In fact, never inline a dllimport function that references a non-imported function or variable (because these are not defined in the importing library). This seems to match MSVC's behaviour. Differential Revision: http://reviews.llvm.org/D12422 llvm-svn: 246338
* [CUDA] Change initializer for CUDA device code based on CUDA documentation.Jingyue Wu2015-08-221-1/+21
| | | | | | | | | | | | | | | | | | | | | | | | | | Summary: According to CUDA documentation, global variables declared with __device__, __constant__ can be initialized from host code, so mark them as externally initialized. Because __shared__ variables cannot have an initialization as part of their declaration and since the value maybe kept across different kernel invocation, the value of __shared__ is effectively undefined instead of zero initialized. Wrongly using zero initializer may cause illegitimate optimization, e.g. removing unused __constant__ variable because it's not updated in the device code and the value is initialized with zero. Test Plan: test/CodeGenCUDA/address-spaces.cu Patch by Xuetian Weng Reviewers: jholewinski, eliben, tra, jingyue Subscribers: llvm-commits Differential Revision: http://reviews.llvm.org/D12241 llvm-svn: 245786
* [modules] Don't eagerly deserialize so many ImportDecls. CodeGen basically ↵Richard Smith2015-08-191-5/+2
| | | | | | ignores ImportDecls imported from modules, so only eagerly deserialize the ones from a PCH / preamble. llvm-svn: 245406
OpenPOWER on IntegriCloud