llvm-project/clang/test/CodeGenCUDA
Justin Lebar e56360a2cd [CUDA] Align kernel launch args correctly when the LLVM type's alignment is different from the clang type's alignment.
Summary:
Before this patch, we computed the offsets in memory of args passed to
GPU kernel functions by throwing all of the args into an LLVM struct.

clang emits packed llvm structs basically whenever it feels like it, and
packed structs have alignment 1.  So we cannot rely on the llvm type's
alignment matching the C++ type's alignment.

This patch fixes our codegen so we always respect the clang types'
alignments.

Reviewers: rnk

Subscribers: cfe-commits, tra

Differential Revision: https://reviews.llvm.org/D22879

llvm-svn: 276927
2016-07-27 22:36:21 +00:00
..
Inputs Avoid depending on test inputes that aren't in Inputs 2016-05-20 00:38:25 +00:00
address-spaces.cu [CUDA] Restrict init of local __shared__ variables to empty constructors only. 2016-05-09 22:09:56 +00:00
alias.cu [CUDA] Don't generate aliases for static extern "C" functions. 2016-01-25 22:36:37 +00:00
convergent.cu [CUDA] Conservatively mark inline asm as convergent. 2016-05-31 21:27:13 +00:00
cuda-builtin-vars.cu NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx ones 2016-07-07 16:41:08 +00:00
device-stub.cu Fixed test failure platforms with name mangling different from Linux. 2016-03-02 21:03:20 +00:00
device-var-init.cu Avoid depending on test inputes that aren't in Inputs 2016-05-20 00:38:25 +00:00
device-vtable.cu [CUDA] Make vtable construction aware of host/device side of CUDA compilation. 2015-12-17 18:12:36 +00:00
filter-decl.cu [CUDA] Emit host-side 'shadows' for device-side global variables 2016-03-02 18:28:50 +00:00
flush-denormals.cu [CUDA] Fix flush-denormals.cu test so that it checks what it intends to CHECK. 2016-05-10 00:34:50 +00:00
fp-contract.cu [CUDA] Enable fusing FP ops (-ffp-contract=fast) for CUDA by default. 2016-05-19 18:44:45 +00:00
function-overload.cu [CUDA] Remove three obsolete CUDA cc1 flags. 2016-03-29 16:24:16 +00:00
host-device-calls-host.cu [CUDA] Remove three obsolete CUDA cc1 flags. 2016-03-29 16:24:16 +00:00
kernel-args-alignment.cu [CUDA] Align kernel launch args correctly when the LLVM type's alignment is different from the clang type's alignment. 2016-07-27 22:36:21 +00:00
kernel-call.cu
launch-bounds.cu [CUDA] Add implicit conversion of __launch_bounds__ arguments to rvalue. 2016-06-06 22:54:57 +00:00
link-device-bitcode.cu [CUDA] Add -disable-llvm-passes to CodeGenCUDA/link-device-bitcode.cu. NFC 2016-03-30 23:45:38 +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
printf-aggregate.cu [CUDA] Don't crash when trying to printf a non-scalar object. 2016-02-11 02:00:52 +00:00
printf.cu [CUDA] Generate CUDA's printf alloca in its function's entry block. 2016-01-28 23:58:28 +00:00
ptx-kernels.cu [CUDA] Give templated device functions internal linkage, templated kernels external linkage. 2016-06-30 18:41:33 +00:00