summaryrefslogtreecommitdiff
path: root/lib/Sema/SemaCUDA.cpp
Commit message (Collapse)AuthorAgeFilesLines
* [CUDA][HIP] Fix host/device check with -fopenmpYaxun Liu2019-10-091-40/+12
| | | | | | | | | | | | | | | | | CUDA/HIP program may be compiled with -fopenmp. In this case, -fopenmp is only passed to host compilation to take advantages of multi-threads computation. CUDA/HIP and OpenMP both use Sema::DeviceCallGraph to store functions to be analyzed and remove them once they decide the function is sure to be emitted. CUDA/HIP and OpenMP have different functions to determine if a function is sure to be emitted. To check host/device correctly for CUDA/HIP when -fopenmp is enabled, there needs a unified logic to determine whether a function is to be emitted. The logic needs to be aware of both CUDA and OpenMP logic. Differential Revision: https://reviews.llvm.org/D67837 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@374263 91177308-0d34-0410-b5e6-96231b3b80d8
* [HIP] Support new kernel launching APIYaxun Liu2019-09-241-1/+2
| | | | | | | Differential Revision: https://reviews.llvm.org/D67947 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@372773 91177308-0d34-0410-b5e6-96231b3b80d8
* Revert assertion added by r372394Yaxun Liu2019-09-211-1/+0
| | | | | | | | | | The assertion added by r372394 causes CUDA test in test-suite to assert. The assertion was not there originally, so revert it. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@372452 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA][HIP] Fix hostness of defaulted constructorYaxun Liu2019-09-201-12/+27
| | | | | | | | | | | | Clang does not respect the explicit device host attributes of defaulted special members. Also clang does not respect the hostness of special members determined by their first declarations. Clang also adds duplicate implicit device or host attributes in certain cases. This patch fixes that. Differential Revision: https://reviews.llvm.org/D67509 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@372394 91177308-0d34-0410-b5e6-96231b3b80d8
* Split ActOnCallExpr into an ActOnCallExpr to be called by the parser,Richard Smith2019-05-081-1/+1
| | | | | | | and a BuildCallExpr to be called internally within Sema to build / rebuild calls. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@360217 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA][HIP][Sema] Fix template kernel with function as template parameterYaxun Liu2019-03-051-0/+5
| | | | | | | | | | | | | | | | | | | | | If a kernel template has a function as its template parameter, a device function should be allowed as template argument since a kernel can call a device function. However, currently if the kernel template is instantiated in a host function, clang will emit an error message saying the device function is an invalid candidate for the template parameter. This happens because clang checks the reference to the device function during parsing the template arguments. At this point, the template is not instantiated yet. Clang incorrectly assumes the device function is called by the host function and emits the error message. This patch fixes the issue by disabling checking of device function during parsing template arguments and deferring the check to the instantion of the template. At that point, the template decl is already available, therefore the check can be done against the instantiated function template decl. Differential Revision: https://reviews.llvm.org/D56411 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@355421 91177308-0d34-0410-b5e6-96231b3b80d8
* [SEMA]Generalize deferred diagnostic interface, NFC.Alexey Bataev2019-02-071-194/+35
| | | | | | | | | | | | | | | | | Summary: Deferred diagnostic interface is going to be used for OpenMP device compilation. Generalized previously existed deferred diagnostic interface for CUDA to be used with OpenMP and, possibly, other models. Reviewers: rjmccall, tra Subscribers: caomhin, cfe-commits, kkwli0 Tags: #clang Differential Revision: https://reviews.llvm.org/D57908 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@353456 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] add support for the new kernel launch API in CUDA-9.2+.Artem Belevich2019-01-311-3/+16
| | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@352799 91177308-0d34-0410-b5e6-96231b3b80d8
* Update the file headers across all of the LLVM projects in the monorepoChandler Carruth2019-01-191-4/+3
| | | | | | | | | | | | | | | | | to reflect the new license. We understand that people may be surprised that we're moving the header entirely to discuss the new license. We checked this carefully with the Foundation's lawyer and we believe this is the correct approach. Essentially, all code in the project is now made available by the LLVM project under our new license, so you will see that the license headers include that license only. Some of our contributors have contributed code under our old license, and accordingly, we have retained a copy of our old license notice in the top-level files in each project and repository. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@351636 91177308-0d34-0410-b5e6-96231b3b80d8
* [AST][NFC] Pass the AST context to one of the ctor of DeclRefExpr.Bruno Ricci2018-12-211-1/+1
| | | | | | | | | | | All of the other constructors already take a reference to the AST context. This avoids calling Decl::getASTContext in most cases. Additionally move the definition of the constructor from Expr.h to Expr.cpp since it is calling DeclRefExpr::computeDependence. NFC. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@349901 91177308-0d34-0410-b5e6-96231b3b80d8
* [NFC] Rename clang::AttributeList to clang::ParsedAttrErich Keane2018-07-131-5/+5
| | | | | | | | Since The type no longer contains the 'next' item anymore, it isn't a list, so rename it to ParsedAttr to be more accurate. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@337005 91177308-0d34-0410-b5e6-96231b3b80d8
* AttributeList de-listifying:Erich Keane2018-07-121-4/+5
| | | | | | | | | | | | | Basically, "AttributeList" loses all list-like mechanisms, ParsedAttributes is switched to use a TinyPtrVector (and a ParsedAttributesView is created to have a non-allocating attributes list). DeclaratorChunk gets the later kind, Declarator/DeclSpec keep ParsedAttributes. Iterators are added to the ParsedAttribute types so that for-loops work. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@336945 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Check initializers of instantiated template variables.Artem Belevich2018-06-061-0/+53
| | | | | | | | | We were already performing checks on non-template variables, but the checks on templated ones were missing. Differential Revision: https://reviews.llvm.org/D45231 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@334143 91177308-0d34-0410-b5e6-96231b3b80d8
* Remove \brief commands from doxygen comments.Adrian Prantl2018-05-091-1/+1
| | | | | | | | | | | | | | | | | | | This is similar to the LLVM change https://reviews.llvm.org/D46290. We've been running doxygen with the autobrief option for a couple of years now. This makes the \brief markers into our comments redundant. Since they are a visual distraction and we don't want to encourage more \brief markers in new code either, this patch removes them all. Patch produced by for i in $(git grep -l '\@brief'); do perl -pi -e 's/\@brief //g' $i & done for i in $(git grep -l '\\brief'); do perl -pi -e 's/\\brief //g' $i & done Differential Revision: https://reviews.llvm.org/D46320 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@331834 91177308-0d34-0410-b5e6-96231b3b80d8
* [HIP] Add hip input kind and codegen for kernel launchingYaxun Liu2018-04-251-2/+3
| | | | | | | | | | | | | | | | | | | | | | | | HIP is a language similar to CUDA (https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md ). The language syntax is very similar, which allows a hip program to be compiled as a CUDA program by Clang. The main difference is the host API. HIP has a set of vendor neutral host API which can be implemented on different platforms. Currently there is open source implementation of HIP runtime on amdgpu target (https://github.com/ROCm-Developer-Tools/HIP). This patch adds support of input kind and language standard hip. When hip file is compiled, both LangOpts.CUDA and LangOpts.HIP is turned on. This allows compilation of hip program as CUDA in most cases and only special handling of hip program is needed LangOpts.HIP is checked. This patch also adds support of kernel launching of HIP program using HIP host API. When -x hip is not specified, there is no behaviour change for CUDA. Patch by Greg Rodgers. Revised and lit test added by Yaxun Liu. Differential Revision: https://reviews.llvm.org/D44984 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@330790 91177308-0d34-0410-b5e6-96231b3b80d8
* Revert "[CUDA] Check initializers of instantiated template variables."Artem Belevich2018-04-041-53/+0
| | | | | | | This (temporarily) reverts commit r329127 due to the problems it exposed in TensorFlow. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@329229 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Check initializers of instantiated template variables.Artem Belevich2018-04-031-0/+53
| | | | | | | | | We were already performing checks on non-template variables, but the checks on templated ones were missing. Differential Revision: https://reviews.llvm.org/D45231 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@329127 91177308-0d34-0410-b5e6-96231b3b80d8
* Fix some handling of AST nodes with diagnostics.Richard Trieu2018-03-281-1/+1
| | | | | | | | | | | 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. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@328688 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Fixed false error reporting in case of calling H->G->HD->D.Artem Belevich2018-03-231-3/+6
| | | | | | | | | | Launching a kernel from the host code does not generate code for the kernel itself. This fixes an issue with clang erroneously reporting an error for a HD->D call from within the kernel. Differential Revision: https://reviews.llvm.org/D44837 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@328362 91177308-0d34-0410-b5e6-96231b3b80d8
* Function with unparsed body is a definitionSerge Pavlov2017-06-211-6/+0
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | While a function body is being parsed, the function declaration is not considered as a definition because it does not have a body yet. In some cases it leads to incorrect interpretation, the case is presented in https://bugs.llvm.org/show_bug.cgi?id=14785: ``` template<typename T> struct Somewhat { void internal() const {} friend void operator+(int const &, Somewhat<T> const &) {} }; void operator+(int const &, Somewhat<char> const &x) { x.internal(); } ``` When statement `x.internal()` in the body of global `operator+` is parsed, the type of `x` must be completed, so the instantiation of `Somewhat<char>` is started. It instantiates the declaration of `operator+` defined inline, and makes a check for redefinition. The check does not detect another definition because the declaration of `operator+` is still not defining as does not have a body yet. To solves this problem the function `isThisDeclarationADefinition` considers a function declaration as a definition if it has flag `WillHaveBody` set. This change fixes PR14785. Differential Revision: https://reviews.llvm.org/D30375 This is a recommit of 305379, reverted in 305381, with small changes. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@305903 91177308-0d34-0410-b5e6-96231b3b80d8
* Reverted 305379 (Function with unparsed body is a definition)Serge Pavlov2017-06-141-0/+6
| | | | | | | | It broke clang-x86_64-linux-selfhost-modules-2 and some other buildbots. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@305381 91177308-0d34-0410-b5e6-96231b3b80d8
* Function with unparsed body is a definitionSerge Pavlov2017-06-141-6/+0
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | While a function body is being parsed, the function declaration is not considered as a definition because it does not have a body yet. In some cases it leads to incorrect interpretation, the case is presented in https://bugs.llvm.org/show_bug.cgi?id=14785: ``` template<typename T> struct Somewhat { void internal() const {} friend void operator+(int const &, Somewhat<T> const &) {} }; void operator+(int const &, Somewhat<char> const &x) { x.internal(); } ``` When statement `x.internal()` in the body of global `operator+` is parsed, the type of `x` must be completed, so the instantiation of `Somewhat<char>` is started. It instantiates the declaration of `operator+` defined inline, and makes a check for redefinition. The check does not detect another definition because the declaration of `operator+` is still not defining as does not have a body yet. To solves this problem the function `isThisDeclarationADefinition` considers a function declaration as a definition if it has flag `WillHaveBody` set. This change fixes PR14785. Differential Revision: https://reviews.llvm.org/D30375 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@305379 91177308-0d34-0410-b5e6-96231b3b80d8
* Factor out some common code between SpecialMemberExceptionSpecInfo and ↵Richard Smith2017-02-241-8/+6
| | | | | | | | | SpecialMemberDeletionInfo. To simplify this, convert SpecialMemberOverloadResult to a value type. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@296073 91177308-0d34-0410-b5e6-96231b3b80d8
* [Sema] Replace remove_if+erase with erase_if. NFC.George Burgess IV2017-01-041-4/+2
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@290991 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Ignore implicit target attributes during function template instantiation.Artem Belevich2016-12-081-6/+33
| | | | | | | | | | | | | | | | | | | | | | | Some functions and templates are treated as __host__ __device__ even when they don't have explicitly specified target attributes. What's worse, this treatment may change depending on command line options (-fno-cuda-host-device-constexpr) or #pragma clang force_cuda_host_device. Combined with strict checking for matching function target that comes with D25809(r288962), it makes it hard to write code which would explicitly instantiate or specialize some functions regardless of pragmas or command line options in effect. This patch changes the way we match target attributes of base template vs attributes used in explicit instantiation or specialization so that only explicitly specified attributes are considered. This makes base template selection behave consistently regardless of pragma of command line options that may affect CUDA target. Differential Revision: https://reviews.llvm.org/D25845 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@289091 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Improve target attribute checking for function templates.Artem Belevich2016-12-071-0/+68
| | | | | | | | | | | | * __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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@288962 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Use only the GVALinkage on function definitions.Justin Lebar2016-11-081-2/+16
| | | | | | | | | | | | | | | | | | | Summary: Previously we'd look at the GVALinkage of whatever FunctionDecl you happened to be calling. This is not right. In the absence of the gnu_inline attribute, to be handled separately, the function definition determines the function's linkage. So we need to wait until we get a def before we can know whether something is known-emitted. Reviewers: tra Subscribers: cfe-commits, rsmith Differential Revision: https://reviews.llvm.org/D26268 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@286313 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Use FunctionDeclAndLoc for the Sema::LocsWithCUDACallDiags hashtable.Justin Lebar2016-10-211-1/+1
| | | | | | | | | | | | Summary: NFC Reviewers: rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25797 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284869 91177308-0d34-0410-b5e6-96231b3b80d8
* Removed unused function argument. NFC.Artem Belevich2016-10-211-1/+1
| | | | | | Differential Revision: https://reviews.llvm.org/D25839 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284843 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] When we emit an error that might have been deferred, also print a ↵Justin Lebar2016-10-191-44/+98
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284647 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Emit errors for wrong-side calls made on the same line as ↵Justin Lebar2016-10-191-4/+6
| | | | | | | | | | | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284643 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Fix false-positive in known-emitted handling.Justin Lebar2016-10-171-6/+22
| | | | | | | | | | | | | | | | | 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. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284355 91177308-0d34-0410-b5e6-96231b3b80d8
* Add and use isDiscardableGVALinkage function.Justin Lebar2016-10-131-1/+1
| | | | | | | | | | Reviewers: rnk Subscribers: cfe-commits Differential Revision: https://reviews.llvm.org/D25571 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284159 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Emit deferred diagnostics during Sema rather than during codegen.Justin Lebar2016-10-131-38/+173
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284158 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Add Sema::CUDADiagBuilder and Sema::CUDADiagIf{Device,Host}Code().Justin Lebar2016-10-131-72/+83
| | | | | | | | | | | | | | | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@284143 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Make touching a kernel from a __host__ __device__ function a deferred ↵Justin Lebar2016-10-121-2/+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. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283963 91177308-0d34-0410-b5e6-96231b3b80d8
* Aligned allocation versus CUDA: make deallocation function preference orderRichard Smith2016-10-111-62/+11
| | | | | | | | | | | | | | 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. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283830 91177308-0d34-0410-b5e6-96231b3b80d8
* Re-commit r283722, reverted in r283750, with a fix for a CUDA-specific use ofRichard Smith2016-10-101-0/+28
| | | | | | | | | | | | past-the-end iterator. Original commit message: P0035R4: Semantic analysis and code generation for C++17 overaligned allocation. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283789 91177308-0d34-0410-b5e6-96231b3b80d8
* Revert "P0035R4: Semantic analysis and code generation for C++17 overaligned ↵Daniel Jasper2016-10-101-28/+0
| | | | | | | | | | | | allocation." This reverts commit r283722. Breaks: Clang.SemaCUDA.device-var-init.cu Clang.CodeGenCUDA.device-var-init.cu http://lab.llvm.org:8080/green/job/clang-stage1-cmake-RA-expensive/884/ git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283750 91177308-0d34-0410-b5e6-96231b3b80d8
* P0035R4: Semantic analysis and code generation for C++17 overalignedRichard Smith2016-10-101-0/+28
| | | | | | | allocation. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283722 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Add #pragma clang force_cuda_host_device_{begin,end} pragmas.Justin Lebar2016-10-081-0/+27
| | | | | | | | | | | | | | | | | Summary: These cause us to consider all functions in-between to be __host__ __device__. You can nest these pragmas; you just can't have more 'end's than 'begin's. Reviewers: rsmith Subscribers: tra, jhen, cfe-commits Differential Revision: https://reviews.llvm.org/D24975 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283677 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Do a better job at detecting wrong-side calls.Justin Lebar2016-10-081-1/+7
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283637 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Harmonize asserts in SemaCUDA, NFC.Justin Lebar2016-09-301-3/+3
| | | | git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282987 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Remove incorrect comment in CUDASetLambdaAttrs.Justin Lebar2016-09-301-4/+0
| | | | | | | | | I'd said that nvcc doesn't allow you to add __host__ or __device__ attributes on lambdas in all circumstances, but I believe this was user error on my part. I can't reproduce these warnings/errors if I pass --expt-extended-lambda to nvcc. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282912 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Make lambdas inherit __host__ and __device__ attributes from the ↵Justin Lebar2016-09-301-0/+19
| | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282880 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Disallow variable-length arrays in CUDA device code.Justin Lebar2016-09-281-0/+20
| | | | | | | | | | Reviewers: tra Subscribers: cfe-commits, jhen Differential Revision: https://reviews.llvm.org/D25050 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282647 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Disallow exceptions in device code.Justin Lebar2016-09-281-0/+24
| | | | | | | | | | Reviewers: tra Subscribers: cfe-commits, jhen Differential Revision: https://reviews.llvm.org/D25036 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@282646 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Fix "declared here" note on deferred wrong-side errors.Justin Lebar2016-08-161-5/+10
| | | | | | | Previously we weren't deferring these "declared here" notes, which is obviously wrong. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@278767 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Raise an error if a wrong-side call is codegen'ed.Justin Lebar2016-08-151-0/+30
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@278759 91177308-0d34-0410-b5e6-96231b3b80d8
* [CUDA] Use the multi-element remove function in EraseUnwantedCUDAMatches.Justin Lebar2016-07-121-2/+4
| | | | | | | | | | | | | | | | Summary: Bug pointed out by Benjamin Kramer in r264008. I think the bug is benign because by the time this is called, we should only have at most two overloads to consider (either a host and a device overload, or a host+device overload, but not all three). Reviewers: tra Subscribers: cfe-commits, bkramer Differential Revision: http://reviews.llvm.org/D21914 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@275233 91177308-0d34-0410-b5e6-96231b3b80d8
View Lorry Controller Status