summaryrefslogtreecommitdiffstats
path: root/clang/test/SemaCUDA/function-overload.cu
Commit message (Collapse)AuthorAgeFilesLines
* [CUDA][HIP] Re-apply part of r372318.Michael Liao2019-09-191-0/+17
| | | | | | | | | | - r372318 causes violation of `use-of-uninitialized-value` detected by MemorySanitizer. Once `Viable` field is set to false, `FailureKind` needs setting as well as it will be checked during destruction if `Viable` is not true. - Revert the part trying to skip `std::vector` erasing. llvm-svn: 372356
* Revert "[CUDA][HIP] Fix typo in `BestViableFunction`"Mitch Phillips2019-09-191-17/+0
| | | | | | | | Broke the msan buildbots (see comments on rL372318 for more details). This reverts commit eb231d15825ac345b546f4c99372d1cac8f14f02. llvm-svn: 372353
* [CUDA][HIP] Fix typo in `BestViableFunction`Michael Liao2019-09-191-0/+17
| | | | | | | | | | | | | | | | | | | Summary: - Should consider viable ones only when checking SameSide candidates. - Replace erasing with clearing viable flag to reduce data moving/copying. - Add one and revise another one as the diagnostic message are more relevant compared to previous one. Reviewers: tra Subscribers: cfe-commits, yaxunl Tags: #clang Differential Revision: https://reviews.llvm.org/D67730 llvm-svn: 372318
* Fix some handling of AST nodes with diagnostics.Richard Trieu2018-03-281-2/+2
| | | | | | | | | The diagnostic system for Clang can already handle many AST nodes. Instead of converting them to strings first, just hand the AST node directly to the diagnostic system and let it handle the output. Minor changes in some diagnostic output. llvm-svn: 328688
* PR34163: Don't cache an incorrect key function for a class if queried betweenRichard Smith2017-08-121-1/+1
| | | | | | | | | | | the class becoming complete and its inline methods being parsed. This replaces the hack of using the "late parsed template" flag to track member functions with bodies we've not parsed yet; instead we now use the "will have body" flag, which carries the desired implication that the function declaration *is* a definition, and that we've just not parsed its body yet. llvm-svn: 310776
* [CUDA] Improve target attribute checking for function templates.Artem Belevich2016-12-071-26/+32
| | | | | | | | | | | | * __host__ __device__ functions are no longer considered to be redeclarations of __host__ or __device__ functions. This prevents unintentional merging of target attributes across them. * Function target attributes are not considered (and must match) during explicit instantiation and specialization of function templates. Differential Revision: https://reviews.llvm.org/D25809 llvm-svn: 288962
* [CUDA] Rework tests now that we emit deferred diagnostics during sema. ↵Justin Lebar2016-10-191-0/+10
| | | | | | | | | | | | | | | | | | | | | | Test-only change. Summary: Previously we had to split out a lot of our tests into a test that checked only immediate errors and a test that checked only deferred errors. This was because, if you emitted any immediate errors, we wouldn't run codegen, where the deferred errors were emitted. We've fixed this, and now emit deferred errors during sema. This lets us merge a bunch of tests, and lets us convert some other tests to -fsyntax-only. Reviewers: tra Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25755 llvm-svn: 284553
* [CUDA] Emit deferred diagnostics during Sema rather than during codegen.Justin Lebar2016-10-131-2/+42
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: Emitting deferred diagnostics during codegen was a hack. It did work, but usability was poor, both for us as compiler devs and for users. We don't codegen if there are any sema errors, so for users this meant that they wouldn't see deferred errors if there were any non-deferred errors. For devs, this meant that we had to carefully split up our tests so that when we tested deferred errors, we didn't emit any non-deferred errors. This change moves checking for deferred errors into Sema. See the big comment in SemaCUDA.cpp for an overview of the idea. This checking adds overhead to compilation, because we have to maintain a partial call graph. As a result, this change makes deferred errors a CUDA-only concept (whereas before they were a general concept). If anyone else wants to use this framework for something other than CUDA, we can generalize at that time. This patch makes the minimal set of test changes -- after this lands, I'll go back through and do a cleanup of the tests that we no longer have to split up. Reviewers: rnk Subscribers: cfe-commits, rsmith, tra Differential Revision: https://reviews.llvm.org/D25541 llvm-svn: 284158
* [CUDA] Make touching a kernel from a __host__ __device__ function a deferred ↵Justin Lebar2016-10-121-12/+1
| | | | | | | | | | | | | | | | error. Previously, this was an immediate, don't pass go, don't collect $200 error. But this precludes us from writing code like __host__ __device__ void launch_kernel() { kernel<<<...>>>(); } Such code isn't wrong, following our notions of right and wrong in CUDA, unless it's codegen'ed. llvm-svn: 283963
* [CUDA] Disallow overloading destructors.Justin Lebar2016-10-031-33/+0
| | | | | | | | | | | | | | | | | | | | | | Summary: We'd attempted to allow this, but turns out we were doing a very bad job. :) Making this work properly would be a giant change in clang. For example, we'd need to make CXXRecordDecl::getDestructor() context-sensitive, because the destructor you end up with depends on where you're calling it from. For now (and hopefully for ever), just disallow overloading of destructors in CUDA. Reviewers: rsmith Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D24571 llvm-svn: 283120
* [CUDA] Do not merge CUDA target attributes.Artem Belevich2016-09-131-0/+11
| | | | | | | | | | | | CUDA target attributes are used for function overloading and must not be merged. This fixes a bug where attributes were inherited during function template specialization in CUDA and made it impossible for specialized function to provide its own target attributes. Differential Revision: https://reviews.llvm.org/D24522 llvm-svn: 281406
* [CUDA] Remove three obsolete CUDA cc1 flags.Justin Lebar2016-03-291-34/+4
| | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: * -fcuda-target-overloads Previously unconditionally set to true by the driver. Necessary for correct functioning of the compiler -- our CUDA headers wrapper won't compile without this. * -fcuda-disable-target-call-checks Previously unconditionally set to true by the driver. Necessary to compile almost any external CUDA code -- almost all libraries assume that host+device code can call host or device functions. * -fcuda-allow-host-calls-from-host-device No effect when target overloading is enabled. Reviewers: tra Subscribers: rsmith, cfe-commits Differential Revision: http://reviews.llvm.org/D18416 llvm-svn: 264739
* [CUDA] Merge most of CodeGenCUDA/function-overload.cu into ↵Justin Lebar2016-03-231-111/+215
| | | | | | | | | | | | | | | | | | | | SemaCUDA/function-overload.cu. Summary: Previously we were using the codegen test to ensure that we choose the right overload. But we can do this within sema, with a bit of cleverness. I left the constructor/destructor checks in CodeGen, because these overloads (particularly on the destructors) are hard to check in Sema. Reviewers: tra Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D18386 llvm-svn: 264207
* [CUDA] Simplify SemaCUDA/function-overload.cu test.Justin Lebar2016-03-231-84/+77
| | | | | | | | | | | | | | | | | | Summary: Principally, don't hardcode the line numbers of various notes. This lets us make changes to the test without recomputing linenos everywhere. Instead, just tell -verify that we may get 0 or more notes pointing to the relevant function definitions. Checking that we get exactly the right note isn't so important (and anyway is checked elsewhere). Reviewers: tra Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D18385 llvm-svn: 264206
* [CUDA] do not allow attribute-based overloading for __global__ functions.Artem Belevich2016-02-241-0/+10
| | | | | | | | __global__ functions are present on both host and device side, so providing __host__ or __device__ overloads is not going to do anything useful. llvm-svn: 261778
* [CUDA] Tweak attribute-based overload resolution to match nvcc behavior.Artem Belevich2016-02-121-33/+20
| | | | | | | | | | | | | | | | | | | | | | | | | | This is an artefact of split-mode CUDA compilation that we need to mimic. HD functions are sometimes allowed to call H or D functions. Due to split compilation mode device-side compilation will not see host-only function and thus they will not be considered at all. For clang both H and D variants will become function overloads visible to compiler. Normally target attribute is considered only if C++ rules can not determine which function is better. However in this case we need to ignore functions that would not be present during current compilation phase before we apply normal overload resolution rules. Changes: * introduced another level of call preference to better describe possible call combinations. * removed WrongSide functions from consideration if the set contains SameSide function. * disabled H->D, D->H and G->H calls. These combinations are not allowed by CUDA and we were reluctantly allowing them to work around device-side calls to math functions in std namespace. We no longer need it after r258880. Differential Revision: http://reviews.llvm.org/D16870 llvm-svn: 260697
* [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
OpenPOWER on IntegriCloud