summaryrefslogtreecommitdiffstats
path: root/clang/test/CodeGenCUDA/kernel-args-alignment.cu
Commit message (Collapse)AuthorAgeFilesLines
* LLVM IR: Generate new-style byval-with-Type from ClangTim Northover2019-06-051-1/+1
| | | | | | | | | | | LLVM IR recently added a Type parameter to the byval Attribute, so that when pointers become opaque and no longer have an element type the information will still be present in IR. For now the Type parameter is optional (which is why Clang didn't need this change at the time), but it will become mandatory soon. llvm-svn: 362652
* [CUDA] add support for the new kernel launch API in CUDA-9.2+.Artem Belevich2019-01-311-6/+10
| | | | | | | | | | | | | Instead of calling CUDA runtime to arrange function arguments, the new API constructs arguments in a local array and the kernels are launched with __cudaLaunchKernel(). The old API has been deprecated and is expected to go away in the next CUDA release. Differential Revision: https://reviews.llvm.org/D57488 llvm-svn: 352799
* [CUDA] Align kernel launch args correctly when the LLVM type's alignment is ↵Justin Lebar2016-07-271-0/+36
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
OpenPOWER on IntegriCloud