summaryrefslogtreecommitdiffstats
path: root/clang/lib/Sema/SemaCUDA.cpp
Commit message (Collapse)AuthorAgeFilesLines
...
* Create a frontend flag to disable CUDA cross-target call checksEli Bendersky2015-04-151-0/+5
| | | | | | | | | | | | | | | For CUDA source, Sema checks that the targets of call expressions make sense (e.g. a host function can't call a device function). Adding a flag that lets us skip this check. Motivation: for source-to-source translation tools that have to accept code that's not strictly kosher CUDA but is still accepted by nvcc. The source-to-source translation tool can then fix the code and leave calls that are semantically valid for the actual compilation stage. Differential Revision: http://reviews.llvm.org/D9036 llvm-svn: 235049
* CUDA: Add option to allow host device functions to call host functionsJacques Pienaar2015-02-241-2/+14
| | | | | | Commiting code from review http://reviews.llvm.org/D7841 llvm-svn: 230385
* Consider calls from implict host device functions as valid in SemaCUDA.Jacques Pienaar2014-12-161-6/+6
| | | | | | | | | | | | | | | | | | | | | | | | | | | In SemaCUDA all implicit functions were considered host device, this led to errors such as the following code snippet failing to compile: struct Copyable { const Copyable& operator=(const Copyable& x) { return *this; } }; struct Simple { Copyable b; }; void foo() { Simple a, b; a = b; } Above the implicit copy assignment operator was inferred as host device but there was only a host assignment copy defined which is an error in device compilation mode. Differential Revision: http://reviews.llvm.org/D6565 llvm-svn: 224358
* CUDA host device code with two code pathsReid Kleckner2014-12-031-5/+14
| | | | | | | | | | | | | | | | | | | | | | | Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
* CUDA: mark the target of implicit intrinsics properlyEli Bendersky2014-09-301-0/+6
| | | | | | | | | | | | | r218624 implemented target inference for implicit special members. However, other entities can be implicit - for example intrinsics. These can not have inference running on them, so they should be marked host device as before. This is the safest and most flexible setting, since by construction these functions don't invoke anything, and we'd like them to be invokable from both host and device code. LLVM's intrinsics definitions (where these intrinsics come from in the case of CUDA/NVPTX) have no notion of target, so both host and device intrinsics can be supported this way. llvm-svn: 218688
* CUDA: Fix incorrect target inference for implicit members.Eli Bendersky2014-09-291-4/+176
| | | | | | | | | | | | As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device, etc) for implicit members (constructors, etc.) incorrectly. This causes errors and even assertions in Clang when compiling code (assertions in C++11 mode where implicit move constructors are added into the mix). Fix the problem by inferring the target from the methods the implicit member should call (depending on its base classes and fields). llvm-svn: 218624
* Split off CUDA-specific Sema parts to a new fileEli Bendersky2014-09-031-0/+76
In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
OpenPOWER on IntegriCloud