summaryrefslogtreecommitdiffstats
path: root/clang/test/SemaCUDA
Commit message (Collapse)AuthorAgeFilesLines
...
* [CUDA] Simplify some repeated diagnostic expectations in CUDA tests.Justin Lebar2016-10-214-22/+7
| | | | | | | | Instead of repeating the diagnostic, use "expected-note N". Test-only change. llvm-svn: 284882
* Declare H and H new/delete.Artem Belevich2016-10-211-1/+45
| | | | llvm-svn: 284879
* [CUDA] When we emit an error that might have been deferred, also print a ↵Justin Lebar2016-10-197-4/+51
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | callstack. Summary: Previously, when you did something not allowed in a host+device function and then caused it to be codegen'ed, we would print out an error telling you that you did something bad, but we wouldn't tell you how we decided that the function needed to be codegen'ed. This change causes us to print out a callstack when emitting deferred errors. This is immensely helpful when debugging highly-templated code, where it's often unclear how a function became known-emitted. We only print the callstack once per function, after we print the all deferred errors. This patch also switches all of our hashtables to using canonical FunctionDecls instead of regular FunctionDecls. This prevents a number of bugs, some of which are caught by tests added here, in which we assume that two FDs for the same function have the same pointer value. Reviewers: rnk Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D25704 llvm-svn: 284647
* [CUDA] Emit errors for wrong-side calls made on the same line as ↵Justin Lebar2016-10-191-0/+39
| | | | | | | | | | | | | | | | | | | | | | | non-wrong-side calls. Summary: This fixes two related bugs: 1) Previously, if you had a non-wrong side call at some source code location L, we wouldn't emit errors for wrong-side calls that appeared at L. 2) We'd only emit one wrong-side error per source code location, when we actually want to emit it twice if we hit this line more than once due to e.g. template instantiation. Reviewers: tra Subscribers: rnk, cfe-commits Differential Revision: https://reviews.llvm.org/D25702 llvm-svn: 284643
* [CUDA] Rework tests now that we emit deferred diagnostics during sema. ↵Justin Lebar2016-10-1910-137/+86
| | | | | | | | | | | | | | | | | | | | | | 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] Fix false-positive in known-emitted handling.Justin Lebar2016-10-171-0/+44
| | | | | | | | | | | | | | | | | Previously: When compiling for host, our constructed call graph went *through* kernel calls. This meant that if we had host calls kernel calls HD we would incorrectly mark the HD function as known-emitted on the host side, and thus perform host-side checks on it. Fixing this exposed another issue, wherein when marking a function as known-emitted, we also need to traverse the callgraph of its template, because non-dependent calls are attached to a function's template, not its instantiation. llvm-svn: 284355
* [CUDA] Emit deferred diagnostics during Sema rather than during codegen.Justin Lebar2016-10-134-10/+48
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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] Allow static variables in __host__ __device__ functions, so long as ↵Justin Lebar2016-10-132-3/+23
| | | | | | | | | | | | they're never codegen'ed for device. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25150 llvm-svn: 284145
* [CUDA] Disallow __shared__ variables in host functions.Justin Lebar2016-10-131-0/+1
| | | | | | | | | | Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25143 llvm-svn: 284144
* [CUDA] Add Sema::CUDADiagBuilder and Sema::CUDADiagIf{Device,Host}Code().Justin Lebar2016-10-132-8/+8
| | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: Together these let you easily create diagnostics that - are never emitted for host code - are always emitted for __device__ and __global__ functions, and - are emitted for __host__ __device__ functions iff these functions are codegen'ed. At the moment there are only three diagnostics that need this treatment, but I have more to add, and it's not sustainable to write code for emitting every such diagnostic twice, and from a special wrapper in SemaCUDA.cpp. While we're at it, don't emit the function name in err_cuda_device_exceptions: It's not necessary to print it, and making this work in the new framework in the face of a null value for dyn_cast<FunctionDecl>(CurContext) isn't worth the effort. Reviewers: rnk Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D25139 llvm-svn: 284143
* Added REQUIRED triples to the test that fails on some ARM buildbots.Artem Belevich2016-10-121-0/+3
| | | | llvm-svn: 283964
* [CUDA] Make touching a kernel from a __host__ __device__ function a deferred ↵Justin Lebar2016-10-123-19/+37
| | | | | | | | | | | | | | | | 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
* Aligned allocation versus CUDA: make deallocation function preference orderRichard Smith2016-10-111-0/+13
| | | | | | | | | | | | | match other CUDA preference orders, per discussion with jlebar. We now model this in an attempt to match overload resolution as closely as possible: - First, we throw out all non-callable (due to CUDA host/device mismatch) operator delete functions. - Then we apply sizedness / alignedness preferences based on whether the type is overaligned and whether the deallocation function is a member. - Finally, we use the CUDA callability preference as a tiebreaker. llvm-svn: 283830
* [CUDA] Rename cuda_builtin_vars.h to __clang_cuda_builtin_vars.h.Justin Lebar2016-10-081-6/+6
| | | | | | | | | | | | Summary: This matches the idiom we use for our other CUDA wrapper headers. Reviewers: tra Subscribers: beanz, mgorny, cfe-commits Differential Revision: https://reviews.llvm.org/D24978 llvm-svn: 283679
* [CUDA] Do a better job at detecting wrong-side calls.Justin Lebar2016-10-082-1/+42
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: Move CheckCUDACall from ActOnCallExpr and BuildDeclRefExpr to DiagnoseUseOfDecl. This lets us catch some edge cases we were missing, specifically around class operators. This necessitates a few other changes: - Avoid emitting duplicate deferred diags in CheckCUDACall. Previously we'd carefully placed our call to CheckCUDACall such that it would only ever run once for a particular callsite. But now this isn't the case. - Emit deferred diagnostics from a template specialization/instantiation's primary template, in addition to from the specialization/instantiation itself. DiagnoseUseOfDecl ends up putting the deferred diagnostics on the template, rather than the specialization, so we need to check both. Reviewers: rsmith Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D24573 llvm-svn: 283637
* [CUDA] Disallow overloading destructors.Justin Lebar2016-10-033-50/+33
| | | | | | | | | | | | | | | | | | | | | | 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] Allow extern __shared__ on empty-length arrays.Justin Lebar2016-10-021-1/+10
| | | | | | "extern __shared__ int x[]" is OK. llvm-svn: 283068
* [CUDA] Disallow __constant__ local variables.Justin Lebar2016-09-301-0/+8
| | | | | | | | | | Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25129 llvm-svn: 282986
* [CUDA] Disallow 'extern __shared__' variables.Justin Lebar2016-09-302-0/+16
| | | | | | | | | | | | | | | | | | Also add a test that we disallow __constant__ __shared__ int x; because it's possible to break this without breaking __shared__ __constant__ int x; Reviewers: rnk Subscribers: cfe-commits, tra Differential Revision: https://reviews.llvm.org/D25125 llvm-svn: 282985
* [CUDA] Fix implicit-device-lambda.cu after r282911.Justin Lebar2016-09-301-2/+2
| | | | | | | This commit added a warning that we're (correctly) hitting in this test. Just ignore it. llvm-svn: 282927
* [CUDA] Make lambdas inherit __host__ and __device__ attributes from the ↵Justin Lebar2016-09-302-0/+113
| | | | | | | | | | | | | | scope in which they're created. Summary: NVCC compat. Fixes bug 30567. Reviewers: tra Subscribers: cfe-commits, rnk Differential Revision: https://reviews.llvm.org/D25105 llvm-svn: 282880
* [CUDA] Disallow variable-length arrays in CUDA device code.Justin Lebar2016-09-282-0/+33
| | | | | | | | | | Reviewers: tra Subscribers: cfe-commits, jhen Differential Revision: https://reviews.llvm.org/D25050 llvm-svn: 282647
* [CUDA] Disallow exceptions in device code.Justin Lebar2016-09-282-0/+59
| | | | | | | | | | Reviewers: tra Subscribers: cfe-commits, jhen Differential Revision: https://reviews.llvm.org/D25036 llvm-svn: 282646
* [AMDGPU] Expose flat work group size, register and wave control attributesKonstantin Zhuravlyov2016-09-262-14/+110
| | | | | | | | | __attribute__((amdgpu_flat_work_group_size(<min>, <max>))) - request minimum and maximum flat work group size __attribute__((amdgpu_waves_per_eu(<min>[, <max>]))) - request minimum and/or maximum waves per execution unit Differential Revision: https://reviews.llvm.org/D24513 llvm-svn: 282371
* [CUDA] Add test checking our ability to take a function pointer to a ↵Justin Lebar2016-09-141-0/+31
| | | | | | | | | | | | | | __global__ function on the host side. Summary: This functionality is used by Thrust. Reviewers: tra Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D24581 llvm-svn: 281543
* [CUDA] Do not merge CUDA target attributes.Artem Belevich2016-09-132-0/+40
| | | | | | | | | | | | 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] Fix "declared here" note on deferred wrong-side errors.Justin Lebar2016-08-162-0/+19
| | | | | | | Previously we weren't deferring these "declared here" notes, which is obviously wrong. llvm-svn: 278767
* [CUDA] Raise an error if a wrong-side call is codegen'ed.Justin Lebar2016-08-153-0/+169
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: Some function calls in CUDA are allowed to appear in semantically-correct programs but are an error if they're ever codegen'ed. Specifically, a host+device function may call a host function, but it's an error if such a function is ever codegen'ed in device mode (and vice versa). Previously, clang made no attempt to catch these errors. For the most part, they would be caught by ptxas, and reported as "call to unknown function 'foo'". Now we catch these errors and report them the same as we report other illegal calls (e.g. a call from a host function to a device function). This has a small change in error-message behavior for calls that were previously disallowed (e.g. calls from a host to a device function). Previously, we'd catch disallowed calls fairly early, before doing additional semantic checking e.g. of the call's arguments. Now we catch these illegal calls at the very end of our semantic checks, so we'll only emit a "illegal CUDA call" error if the call is otherwise well-formed. Reviewers: tra, rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23242 llvm-svn: 278759
* [CUDA] Reject calls to __device__ functions from host variable global ↵Justin Lebar2016-08-101-0/+32
| | | | | | | | | | | | initializers. Reviewers: tra Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23335 llvm-svn: 278196
* [CUDA] Print a "previous-decl" note when calling an illegal member fn.Justin Lebar2016-08-101-3/+3
| | | | | | | | | | | | | | | Summary: When we emit err_ref_bad_target, we should emit a "'method' declared here" note. We already do so in most places, just not in BuildCallToMemberFunction. Reviewers: tra Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D23240 llvm-svn: 278195
* [CUDA] Add additional testcases for EraseUnwantedCUDAMatches.Justin Lebar2016-07-122-0/+49
| | | | | | | | | | | | | | | | | Summary: Specifically, this patch adds testcases for all three calls to EraseUnwantedCUDAMatches. The addr-of-overloaded-fn test I accidentally neutered in r264207, which moved much of CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu. The coverage from overloaded-delete test is new. Reviewers: tra Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D21913 llvm-svn: 275232
* [CUDA] Don't assume that destructors can't be overloaded.Justin Lebar2016-07-121-0/+17
| | | | | | | | | | | | | | | Summary: You can overload a destructor in CUDA, and SemaOverload needs to be tweaked not to crash when it sees an explicit call to an overloaded destructor. Reviewers: rsmith Subscribers: cfe-commits, tra Differential Revision: http://reviews.llvm.org/D21912 llvm-svn: 275231
* NVPTX: Use the nvvm builtins to read SRegs rather than the legacy ptx onesJustin Bogner2016-07-071-3/+3
| | | | | | The ptx spellings were removed from LLVM in r274769. llvm-svn: 274770
* [CUDA] Add implicit conversion of __launch_bounds__ arguments to rvalue.Artem Belevich2016-06-061-0/+6
| | | | | | | | Fixes clang crash reported in PR27778. Differential Revision: http://reviews.llvm.org/D20985 llvm-svn: 271951
* Avoid depending on test inputes that aren't in InputsReid Kleckner2016-05-202-5/+152
| | | | | | | | | | Some people have weird CI systems that run each test subdirectory independently without access to other parallel trees. Unfortunately, this means we have to suffer some duplication until Art can sort out how to share these types. llvm-svn: 270164
* [CUDA] Do not allow non-empty destructors for global device-side variables.Artem Belevich2016-05-191-2/+47
| | | | | | | | | | | | | | According to Cuda Programming guide (v7.5, E2.3.1): > __device__, __constant__ and __shared__ variables defined in namespace > scope, that are of class type, cannot have a non-empty constructor or a > non-empty destructor. Clang already deals with device-side constructors (see D15305). This patch enforces similar rules for destructors. Differential Revision: http://reviews.llvm.org/D20140 llvm-svn: 270108
* [CUDA] Split device-var-init.cu tests into separate Sema and CodeGen parts.Artem Belevich2016-05-191-0/+168
| | | | | | | | | | | Codegen tests for device-side variable initialization are subset of test cases used to verify Sema's part of the job. Including CodeGenCUDA/device-var-init.cu from SemaCUDA makes it easier to keep both sides in sync. Differential Revision: http://reviews.llvm.org/D20139 llvm-svn: 270107
* Add support for derived class special members hiding functions brought in fromRichard Smith2016-05-131-2/+3
| | | | | | | | | a base class via a using-declaration. If a class has a using-declaration declaring either a constructor or an assignment operator, eagerly declare its special members in case they need to displace a shadow declaration from a using-declaration. llvm-svn: 269398
* [CUDA] Make unattributed constexpr functions implicitly host+device.Justin Lebar2016-03-303-0/+97
| | | | | | | | | | | | | | | | | | | | | | | | With this patch, by a constexpr function is implicitly host+device unless: a) it's a variadic function (variadic functions are not allowed on the device side), or b) it's preceeded by a __device__ overload in a system header. The restriction on overloading __host__ __device__ functions on the basis of their CUDA attributes remains in place, but we use (b) to allow us to define __device__ overloads for constexpr functions in cmath, which would otherwise be __host__ __device__ and thus not overloadable. You can disable this behavior with -fno-cuda-host-device-constexpr. Reviewers: tra, rnk, rsmith Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D18380 llvm-svn: 264964
* [CUDA] Remove three obsolete CUDA cc1 flags.Justin Lebar2016-03-297-140/+9
| | | | | | | | | | | | | | | | | | | | | | | | | | | 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] Don't allow templated variadic functions.Justin Lebar2016-03-221-0/+6
| | | | | | | | | | Reviewers: tra Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D18373 llvm-svn: 264106
* [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] Don't specify exact line numbers in cuda-builtin-vars.cu.Justin Lebar2016-02-241-5/+5
| | | | | | | | This makes the test less fragile to changes to cuda_builtin_vars.h. Test-only change. llvm-svn: 261775
* [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] Add -fcuda-allow-variadic-functions.Justin Lebar2016-01-261-6/+15
| | | | | | | | | | | | | | | | | Summary: Turns out the variadic function checking added in r258643 was too strict for some existing users; give them an escape valve. When -fcuda-allow-variadic-functions is passed, the front-end makes no attempt to disallow C-style variadic functions. Calls to va_arg are still not allowed. Reviewers: tra Subscribers: cfe-commits, jhen, echristo, bkramer Differential Revision: http://reviews.llvm.org/D16559 llvm-svn: 258822
* [CUDA] Disallow variadic functions other than printf in device code.Justin Lebar2016-01-231-1/+15
| | | | | | | | | | Reviewers: tra Subscribers: cfe-commits, echristo, jhen Differential Revision: http://reviews.llvm.org/D16484 llvm-svn: 258643
* [CUDA] Reject the alias attribute in CUDA device code.Justin Lebar2016-01-231-0/+11
| | | | | | | | | | | | Summary: CUDA (well, strictly speaking, NVPTX) doesn't support aliases. Reviewers: echristo Subscribers: cfe-commits, jhen, tra Differential Revision: http://reviews.llvm.org/D16502 llvm-svn: 258641
* [CUDA] Bail, rather than crash, on va_arg in device code.Justin Lebar2016-01-201-0/+28
| | | | | | | | | | Reviewers: tra Subscribers: echristo, jhen, cfe-commits Differential Revision: http://reviews.llvm.org/D16331 llvm-svn: 258264
OpenPOWER on IntegriCloud