summaryrefslogtreecommitdiffstats
path: root/clang/test/CodeGenCUDA/address-spaces.cu
Commit message (Collapse)AuthorAgeFilesLines
* [CUDA] Let device-side shared variables be initialized with undefYaxun Liu2018-04-021-9/+10
| | | | | | | | | | | CUDA shared variable should be initialized with undef. Patch by Greg Rodgers. Revised and lit test added by Yaxun Liu. Differential Revision: https://reviews.llvm.org/D44985 llvm-svn: 328994
* [CUDA] Restrict init of local __shared__ variables to empty constructors only.Artem Belevich2016-05-091-31/+0
| | | | | | | | | | Allow only empty constructors for local __shared__ variables in a way identical to restrictions imposed on dynamic initializers for global variables on device. Differential Revision: http://reviews.llvm.org/D20039 llvm-svn: 268982
* [CUDA] Only __shared__ variables can be static local on device side.Artem Belevich2016-05-091-8/+0
| | | | | | | | | | According to CUDA programming guide (v7.5): > E.2.9.4: Within the body of a device or global function, only > shared variables may be declared with static storage class. Differential Revision: http://reviews.llvm.org/D20034 llvm-svn: 268962
* [CUDA] Allow trivial constructors as initializer for __shared__ variables.Artem Belevich2015-09-101-0/+14
| | | | | | Differential Revision: http://reviews.llvm.org/D12739 llvm-svn: 247307
* [CUDA] Change initializer for CUDA device code based on CUDA documentation.Jingyue Wu2015-08-221-3/+3
| | | | | | | | | | | | | | | | | | | | | | | | | | 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
* Fix addrspace when emitting constructors of static local variablesJingyue Wu2015-03-251-0/+17
| | | | | | | | | | | | | | | | | | | | Summary: Due to CUDA's implicit address space casting, the type of a static local variable may be more specific (i.e. with address space qualifiers) than the type expected by the constructor. Emit an addrspacecast in that case. Test Plan: Clang used to crash on the added test. Reviewers: nlewycky, pcc, eliben, rsmith Reviewed By: eliben, rsmith Subscribers: llvm-commits Differential Revision: http://reviews.llvm.org/D8575 llvm-svn: 233208
* Test case updates for explicit type parameter to the gep operatorDavid Blaikie2015-03-131-1/+1
| | | | llvm-svn: 232187
* Update Clang tests to handle explicitly typed load changes in LLVM.David Blaikie2015-02-271-6/+6
| | | | llvm-svn: 230795
* Move all CUDA testing inputs to Inputs/ subdirectory inside the tests.Eli Bendersky2014-04-281-1/+1
| | | | llvm-svn: 207453
* Proper handling of static local variables with address space qualifiers.Eli Bendersky2014-03-241-3/+69
| | | | | | | | Similar to the implementation for globals in r157167. Patch by Jingyue Wu. llvm-svn: 204677
* Fix test failures after addrspacecast added.Matt Arsenault2013-11-151-3/+3
| | | | | | Bitcasts between address spaces are no longer allowed. llvm-svn: 194765
* CUDA: give static storage class to __shared__ and __constant__Peter Collingbourne2012-08-281-0/+8
| | | | | | | | variables without a storage class within a function, to implement CUDA B.2.5: "__shared__ and __constant__ variables have implied static storage [duration]." llvm-svn: 162788
* CUDA: give correct address space to globals declared in functionsPeter Collingbourne2012-08-281-0/+4
| | | | llvm-svn: 162787
* Replace PTX back-end with NVPTX back-end in all places where Clang caresJustin Holewinski2012-05-241-7/+7
| | | | | | NV_CONTRIB llvm-svn: 157403
* CUDA: add CodeGen support for global variable address spaces.Peter Collingbourne2012-05-201-0/+24
Because in CUDA types do not have associated address spaces, globals are declared in their "native" address space, and accessed by bitcasting the pointer to address space 0. This relies on address space 0 being a unified address space. llvm-svn: 157167
OpenPOWER on IntegriCloud