llvm-project/clang/test/CodeGenCUDA
Jingyue Wu 284ebe237f [CUDA] Change initializer for CUDA device code based on CUDA documentation.
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
2015-08-22 05:49:28 +00:00
..
Inputs Move all CUDA testing inputs to Inputs/ subdirectory inside the tests. 2014-04-28 22:21:28 +00:00
address-spaces.cu [CUDA] Change initializer for CUDA device code based on CUDA documentation. 2015-08-22 05:49:28 +00:00
cuda-builtin-vars.cu [cuda] Added support for CUDA built-in variables. 2015-04-21 22:14:13 +00:00
device-stub.cu [cuda] Fixed test case failure on s390x 2015-05-11 18:35:58 +00:00
filter-decl.cu [CUDA] Change initializer for CUDA device code based on CUDA documentation. 2015-08-22 05:49:28 +00:00
host-device-calls-host.cu CUDA: Add option to allow host device functions to call host functions 2015-02-24 21:45:33 +00:00
kernel-call.cu Move all CUDA testing inputs to Inputs/ subdirectory inside the tests. 2014-04-28 22:21:28 +00:00
launch-bounds.cu [cuda] Allow using integral non-type template parameters as launch_bounds attribute arguments. 2015-04-21 22:55:54 +00:00
llvm-used.cu When generating llvm.used, we may need an addrspacecast instead of a bitcast. 2015-02-02 21:05:49 +00:00
ptx-kernels.cu Revert "[CUDA] Add implicit __attribute__((used)) to all __global__ functions." 2015-08-11 11:02:09 +00:00