summaryrefslogtreecommitdiffstats
path: root/clang/test/SemaCUDA
Commit message (Collapse)AuthorAgeFilesLines
...
* [CUDA] Only allow __global__ on free functions and static member functions.Justin Lebar2016-01-201-2/+14
| | | | | | | | | | | | | | Summary: Warn for NVCC compatibility if you declare a static member function or inline function as __global__. Reviewers: tra Subscribers: jhen, echristo, cfe-commits Differential Revision: http://reviews.llvm.org/D16261 llvm-svn: 258263
* [CUDA] Warn undeclared identifiers in CUDA kernel callsJustin Lebar2016-01-142-0/+12
| | | | | | | | | | | | | | | | | | Value, type, and instantiation dependence were not being handled correctly for CUDAKernelCallExpr AST nodes. As a result, if an undeclared identifier was used in the triple-angle-bracket kernel call configuration, there would be no error during parsing, and there would be a crash during code gen. This patch makes sure that an error will be issued during parsing in this case, just as there would be for any other use of an undeclared identifier in C++. Patch by Jason Henline. Reviewers: jlebar, rsmith Differential Revision: http://reviews.llvm.org/D15858 llvm-svn: 257839
* [CUDA] Report an error if code tries to mix incompatible CUDA attributes.Justin Lebar2016-01-133-1/+55
| | | | | | | | | | | | Summary: Thanks to jhen for helping me figure this out. Reviewers: tra, echristo Subscribers: jhen Differential Revision: http://reviews.llvm.org/D16129 llvm-svn: 257554
* Produce a better diagnostic for global register variables.Akira Hatanaka2015-11-181-6/+6
| | | | | | | | | | | | | | | | | | Currently, when there is a global register variable in a program that is bound to an invalid register, clang/llvm prints an error message that is not very user-friendly. This commit improves the diagnostic and moves the check that used to be in the backend to Sema. In addition, it makes changes to error out if the size of the register doesn't match the declared variable size. e.g., volatile register int B asm ("rbp"); rdar://problem/23084219 Differential Revision: http://reviews.llvm.org/D13834 llvm-svn: 253405
* [CUDA] use -aux-triple to pass target triple of opposite side of compilationArtem Belevich2015-11-171-2/+2
| | | | | | | | | | | | | | | | Clang needs to know target triple for both sides of compilation so that preprocessor macros and target builtins from both sides are available. This change augments Compilation class to carry information about toolchains used during different CUDA compilation passes and refactors BuildActions to use it when it constructs CUDA jobs. Removed DeviceTriple from CudaHostAction/CudaDeviceAction as it's no longer needed. Differential Revision: http://reviews.llvm.org/D13144 llvm-svn: 253385
* [CUDA] Allow parsing of host and device code simultaneously.Artem Belevich2015-09-221-19/+14
| | | | | | | | | | | | * adds -aux-triple option to specify target triple * propagates aux target info to AST context and Preprocessor * pulls in target specific preprocessor macros. * pulls in target-specific builtins from aux target. * sets appropriate host or device attribute on builtins. Differential Revision: http://reviews.llvm.org/D12917 llvm-svn: 248299
* [CUDA] Add appropriate host/device attribute to builtins.Artem Belevich2015-09-222-3/+42
| | | | | | | | | The changes are part of attribute-based CUDA function overloading (D12453) and as such are only enabled when it's in effect (-fcuda-target-overloads). Differential Revision: http://reviews.llvm.org/D12122 llvm-svn: 248296
* [CUDA] Allow function overloads in CUDA based on host/device attributes.Artem Belevich2015-09-221-0/+317
| | | | | | | | | | | | | | | | | | | The patch makes it possible to parse CUDA files that contain host/device functions with identical signatures, but different attributes without having to physically split source into host-only and device-only parts. This change is needed in order to parse CUDA header files that have a lot of name clashes with standard include files. Gory details are in design doc here: https://goo.gl/EXnymm Feel free to leave comments there or in this review thread. This feature is controlled with CC1 option -fcuda-target-overloads and is disabled by default. Differential Revision: http://reviews.llvm.org/D12453 llvm-svn: 248295
* [CUDA] Check register names on appropriate side of cuda compilation only.Artem Belevich2015-08-271-4/+28
| | | | | | Differential Revision: http://reviews.llvm.org/D11950 llvm-svn: 246193
* Revert r245496 "[CUDA] Add appropriate host/device attribute to builtins."Artem Belevich2015-08-202-38/+3
| | | | | | It's breaking internal test. llvm-svn: 245592
* [CUDA] Add appropriate host/device attribute to builtins.Artem Belevich2015-08-192-3/+38
| | | | | | Differential Revision: http://reviews.llvm.org/D12122 llvm-svn: 245496
* [CUDA] Added stubs for new attributes used by CUDA headers.Artem Belevich2015-08-101-0/+33
| | | | | | | | | The main purpose is to avoid errors and warnings while parsing CUDA header files. The attributes are currently unused otherwise. Differential version: http://reviews.llvm.org/D11690 llvm-svn: 244497
* [cuda] Preserve TLS storage class of host variable even if it's aArtem Belevich2015-04-281-0/+14
| | | | | | device-side compilation. llvm-svn: 236029
* [cuda] Ignore "TLS unsupported by target" errors for host variables during ↵Artem Belevich2015-04-271-1/+17
| | | | | | | | | | | | | | | | device compilation. During device-side CUDA compilation clang currently complains about all TLS variables, regardless of whether they are __host__ or __device__. This patch suppresses "TLS unsupported" errors for host variables during device compilation and for device variables during host compilation. Differential Revision: http://reviews.llvm.org/D9269 llvm-svn: 235907
* [cuda] Allow using integral non-type template parameters as launch_bounds ↵Artem Belevich2015-04-211-6/+44
| | | | | | | | | | | | | | | | attribute arguments. - Changed CUDALaunchBounds arguments from integers to Expr* so they can be saved in AST for instantiation. - Added support for template instantiation of launch_bounds attrubute. - Moved evaluation of launch_bounds arguments to NVPTXTargetCodeGenInfo:: SetTargetAttributes() where it can be done after template instantiation. - Added a warning on negative launch_bounds arguments. - Amended test cases. Differential Revision: http://reviews.llvm.org/D8985 llvm-svn: 235452
* [cuda] Added support for CUDA built-in variables.Artem Belevich2015-04-211-0/+57
| | | | | | | | | | | | | | | | | Added cuda_builtin_vars.h which implements built-in CUDA variables using __declattr(property). Fields of built-in variables (except for warpSize) are implemented using __declattr(property) which replaces read/write of a member field with a call to a getter/setter member function, in this case with appropriate NVPTX builtin. Added a test case to check diagnostics on attempt to construct or improperly access a built-in variable. Differential Revision: http://reviews.llvm.org/D9064 llvm-svn: 235448
* Revert r235398 "[cuda] Added support for CUDA built-in variables."Artem Belevich2015-04-211-57/+0
| | | | | | r235398 was causing buildbot break due to missing Makefile changes. llvm-svn: 235401
* [cuda] Added support for CUDA built-in variables.Artem Belevich2015-04-211-0/+57
| | | | | | | | | | | | | | | | | Added cuda_builtin_vars.h which implements built-in CUDA variables using __declattr(property). Fields of built-in variables (except for warpSize) are implemented using __declattr(property) which replaces read/write of a member field with a call to a getter/setter member function, in this case with appropriate NVPTX builtin. Added a test case to check diagnostics on attempt to construct or improperly access a built-in variable. Differential Revision: http://reviews.llvm.org/D9064 llvm-svn: 235398
* Create a frontend flag to disable CUDA cross-target call checksEli Bendersky2015-04-151-0/+26
| | | | | | | | | | | | | | | 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
* Ignore device-side asm constraint errors while compiling CUDA code for host ↵Artem Belevich2015-03-192-0/+39
| | | | | | | | and vice versa. Differential Revision: http://reviews.llvm.org/D8392 llvm-svn: 232747
* CUDA: Add option to allow host device functions to call host functionsJacques Pienaar2015-02-242-38/+71
| | | | | | 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-162-4/+57
| | | | | | | | | | | | | | | | | | | | | | | | | | | 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
* Workaround attribute ordering issue with kernel only attributesMatt Arsenault2014-12-051-6/+6
| | | | | | | | | | | Placing the attribute after the kernel keyword would incorrectly reject the attribute, so use the smae workaround that other kernel only attributes use. Also add a FIXME because there are two different phrasings now for the same error, althoug amdgpu_num_[sv]gpr uses a consistent one. llvm-svn: 223490
* Add attributes for AMDGPU register limits.Matt Arsenault2014-12-041-0/+14
| | | | | | | This is a performance hint that can be applied to kernels to attempt to limit the number of used registers. llvm-svn: 223384
* CUDA host device code with two code pathsReid Kleckner2014-12-031-4/+31
| | | | | | | | | | | | | | | | | | | | | | | 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/+10
| | | | | | | | | | | | | 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-293-0/+352
| | | | | | | | | | | | 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
* Fix PR20886 - enforce CUDA target match in method callsEli Bendersky2014-09-251-0/+71
| | | | | | http://reviews.llvm.org/D5298 llvm-svn: 218482
* Automate attribute argument count semantic checking when there are variadic ↵Aaron Ballman2014-07-311-4/+1
| | | | | | | | or optional arguments present. With this, the only time you should have to manually check attribute argument counts is when HasCustomParsing is set to true, or when you have variadic arguments that aren't really variadic (like ownership_holds and friends). Updating the diagnostics in the launch_bounds test since they have been improved in that case. Adding a test for nonnull since it has little test coverage, but has truly variadic arguments. llvm-svn: 214407
* Move all CUDA testing inputs to Inputs/ subdirectory inside the tests.Eli Bendersky2014-04-285-17/+17
| | | | llvm-svn: 207453
* Disallow driver use in more Sema testsAlp Toker2014-04-191-0/+4
| | | | | | | There are now only a handful of Sema tests remaining that use %clang in SemaCXX, SemaObjC and SemaTemplate. llvm-svn: 206688
* Updated the wording of two attribute-related diagnostics so that they print ↵Aaron Ballman2014-01-021-2/+2
| | | | | | the offending attribute name. Also updates the associated test cases. llvm-svn: 198355
* Added a comment about the launch_bounds attribute's AST node being required. ↵Aaron Ballman2013-12-192-0/+15
| | | | | | Since there were no test cases for the attribute, some have been added. This promptly demonstrated a bug with the semantic handling, which is also fixed. llvm-svn: 197637
* CUDA: diagnose invalid calls across targetsPeter Collingbourne2011-10-022-1/+45
| | | | llvm-svn: 140978
* CUDA: add separate diagnostics for too few/many exec config argsPeter Collingbourne2011-10-021-0/+2
| | | | llvm-svn: 140977
* CUDA: diagnose unconfigured calls to global functionsPeter Collingbourne2011-10-021-0/+1
| | | | llvm-svn: 140975
* Sema: diagnose kernel calls to non-global functionsPeter Collingbourne2011-02-231-0/+8
| | | | llvm-svn: 126292
* Parse: add support for parsing CUDA kernel callsPeter Collingbourne2011-02-092-0/+27
| | | | llvm-svn: 125219
* AST, Sema, Serialization: keep track of cudaConfigureCallPeter Collingbourne2011-02-091-0/+3
| | | | llvm-svn: 125216
* Sema: diagnose kernel functions with non-void return typePeter Collingbourne2010-12-121-0/+3
| | | | llvm-svn: 121653
* Basic, Sema: add support for CUDA location attributesPeter Collingbourne2010-12-012-0/+12
llvm-svn: 120545
OpenPOWER on IntegriCloud