Commit Graph

202 Commits

Author SHA1 Message Date
Henry Linjamäki c99b2c6316 CUDA/HIP: Allow __int128 on the host side
Consider case where `__int128` type is supported by the host target but
not by a device target (e.g. spirv*). Clang emits an error message for
unsupported type even if the device code does not use it. This patch
fixes this issue by emitting the error message when the device code
attempts to use the unsupported type.

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D111047
2022-01-04 16:09:26 -08:00
Anshil Gandhi df0560ca00 [HIP] Add atomic load, atomic store and atomic cmpxchng_weak builtin support in HIP-clang
Introduce `__hip_atomic_load`, `__hip_atomic_store` and `__hip_atomic_compare_exchange_weak`
builtins in HIP.

Reviewed By: yaxunl

Differential Revision: https://reviews.llvm.org/D114553
2021-11-29 12:07:13 -07:00
Wei Wang b283d55c90 [openmp] Emit deferred diag only when device compilation presents
There is no need to check for deferred diag when device compilation or target is
not given. This results in considerable build time improvement in some cases.

Differential Revision: https://reviews.llvm.org/D109175
2021-10-25 11:19:18 -07:00
Anshil Gandhi f92db6d3ff [HIP] Relax conditions for address space cast in builtin args
Allow (implicit) address space casting between LLVM-equivalent
target address spaces.

Reviewed By: yaxunl, tra

Differential Revision: https://reviews.llvm.org/D111734
2021-10-15 15:35:52 -06:00
Anshil Gandhi 53fc5100e0 Revert "[HIP] Relax conditions for address space cast in builtin args"
This reverts commit 3b48e1170d.
2021-10-15 14:42:28 -06:00
Anshil Gandhi 3b48e1170d [HIP] Relax conditions for address space cast in builtin args
Allow (implicit) address space casting between LLVM-equivalent
target address spaces.

Reviewed By: yaxunl

Differential Revision: https://reviews.llvm.org/D111734
2021-10-15 14:06:47 -06:00
Artem Belevich 6b20ea6963 [CUDA] Pass ExecConfig through BuildCallToMemberFunction
Otherwise, we fail to compile calls to CUDA kernels that are static members.

Differential Revision: https://reviews.llvm.org/D108787
2021-09-16 11:18:12 -07:00
Yaxun (Sam) Liu 26e492e134 [HIP] Warn capture this pointer in device lambda
HIP currently diagnose capture of this pointer in device lambda in
host member functions. If this pointer points to managed memory,
it can be used in both device and host functions. Under this
situation, capturing this pointer in device lambda functions
in host member functions is valid usage. Change the diagnostic
about capturing this pointer to warning.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D108493
2021-09-08 13:45:26 -04:00
Yaxun (Sam) Liu 82e03e494f [HIP] Defer operator overloading errors
Although clang is able to defer overloading resolution
diagnostics for common functions. It does not defer
overloading resolution caused diagnostics for overloaded
operators.

This patch extends the existing deferred
diagnostic mechanism and defers a diagnostic caused
by overloaded operator.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D104505
2021-06-23 23:39:59 -04:00
Yaxun (Sam) Liu 04caa7c3e0 [CUDA][HIP] Promote const variables to constant
Recently we added diagnosing ODR-use of host variables
in device functions, which includes ODR-use of const
host variables since they are not really emitted on
device side. This caused regressions since we used
to allow ODR-use of const host variables in device
functions.

This patch allows ODR-use of const variables in device
functions if the const variables can be statically initialized
and have an empty dtor. Such variables are marked with
implicit constant attrs and emitted on device side. This is
in line with what clang does for constexpr variables.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D103108
2021-06-01 21:28:41 -04:00
Yaxun (Sam) Liu f7e87dd6ff [CUDA][HIP] Change default lang std to c++14
Currently clang and nvcc use c++14 as default std for C++.
gcc 11 even uses c++17 as default std for C++. However,
clang uses c++98 as default std for CUDA/HIP.

As c++14 has been well adopted and became default for
clang, it seems reasonable to use c++14 as default std
for CUDA/HIP.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D103221
2021-06-01 20:45:10 -04:00
Yaxun (Sam) Liu 4cb42564ec [CUDA][HIP] Fix device variables used by host
variables emitted on both host and device side with different addresses
when ODR-used by host function should not cause device side counter-part
to be force emitted.

This fixes the regression caused by https://reviews.llvm.org/D102237

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D102801
2021-05-20 17:04:29 -04:00
Erich Keane 08ba9ce1ef Suppress Deferred Diagnostics in discarded statements.
It doesn't really make sense to emit language specific diagnostics
in a discarded statement, and suppressing these diagnostics results in a
programming pattern that many users will feel is quite useful.

Basically, this makes sure we only emit errors from the 'true' side of a
'constexpr if'.

It does this by making the ExprEvaluatorBase type have an opt-in option
as to whether it should visit discarded cases.

Differential Revision: https://reviews.llvm.org/D102251
2021-05-12 12:48:47 -07:00
Yaxun (Sam) Liu d8805574c1 [CUDA][HIP] Allow non-ODR use of host var in device
Reviewed by: Artem Belevich, Richard Smith

Differential Revision: https://reviews.llvm.org/D98193
2021-04-19 14:45:24 -04:00
Yaxun (Sam) Liu cc9477166a [CUDA][HIP] add __builtin_get_device_side_mangled_name
Add builtin function __builtin_get_device_side_mangled_name
to get device side manged name for functions and global
variables, which can be used to get symbol address of kernels
or variables by mangled name in dynamically loaded
bundled code objects at run time.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D99301
2021-03-25 15:25:29 -04:00
Yaxun (Sam) Liu 47acdec1dd [CUDA][HIP] Support accessing static device variable in host code for -fgpu-rdc
For -fgpu-rdc mode, static device vars in different TU's may have the same name.
To support accessing file-scope static device variables in host code, we need to give them
a distinct name and external linkage. This can be done by postfixing each static device variable with
a distinct CUID (Compilation Unit ID) hash.

Since the static device variables have different name across compilation units, now we let
them have external linkage so that they can be looked up by the runtime.

Reviewed by: Artem Belevich, and Jon Chesterfield

Differential Revision: https://reviews.llvm.org/D85223
2021-02-24 18:23:45 -05:00
Yaxun (Sam) Liu e355110040 [CUDA][HIP] Fix checking dependent initalizer
Defer constant checking of dependent initializer to template instantiation
since it cannot be done for dependent values.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D95840
2021-02-04 18:04:54 -05:00
Yaxun (Sam) Liu 622eaa4a4c [HIP] Support __managed__ attribute
This patch implements codegen for __managed__ variable attribute for HIP.

Diagnostics will be added later.

Differential Revision: https://reviews.llvm.org/D94814
2021-01-22 11:43:58 -05:00
Artem Belevich 127091bfd5 [CUDA] Normalize handling of defauled dtor.
Defaulted destructor was treated inconsistently, compared to other
compiler-generated functions.

When Sema::IdentifyCUDATarget() got called on just-created dtor which didn't
have implicit __host__ __device__ attributes applied yet, it would treat it as a
host function.  That happened to (sometimes) hide the error when dtor referred
to a host-only functions.

Even when we had identified defaulted dtor as a HD function, we still treated it
inconsistently during selection of usual deallocators, where we did not allow
referring to wrong-side functions, while it is allowed for other HD functions.

This change brings handling of defaulted dtors in line with other HD functions.

Differential Revision: https://reviews.llvm.org/D94732
2021-01-21 10:48:07 -08:00
Artem Belevich 0936655bac [CUDA] Do not diagnose host/device variable access in dependent types.
`isCUDADeviceBuiltinSurfaceType()`/`isCUDADeviceBuiltinTextureType()` do not
work on dependent types as they rely on specific type attributes.

Differential Revision: https://reviews.llvm.org/D92893
2020-12-14 11:53:18 -08:00
Yaxun (Sam) Liu acb6f80d96 [CUDA][HIP] Fix overloading resolution
This patch implements correct hostness based overloading resolution
in isBetterOverloadCandidate.

Based on hostness, if one candidate is emittable whereas the other
candidate is not emittable, the emittable candidate is better.

If both candidates are emittable, or neither is emittable based on hostness, then
other rules should be used to determine which is better. This is because
hostness based overloading resolution is mostly for determining
viability of a function. If two functions are both viable, other factors
should take precedence in preference.

If other rules cannot determine which is better, CUDA preference will be
used again to determine which is better.

However, correct hostness based overloading resolution
requires overloading resolution diagnostics to be deferred,
which is not on by default. The rationale is that deferring
overloading resolution diagnostics may hide overloading reslolutions
issues in header files.

An option -fgpu-exclude-wrong-side-overloads is added, which is off by
default.

When -fgpu-exclude-wrong-side-overloads is off, keep the original behavior,
that is, exclude wrong side overloads only if there are same side overloads.
This may result in incorrect overloading resolution when there are no
same side candates, but is sufficient for most CUDA/HIP applications.

When -fgpu-exclude-wrong-side-overloads is on, enable deferring
overloading resolution diagnostics and enable correct hostness
based overloading resolution, i.e., always exclude wrong side overloads.

Differential Revision: https://reviews.llvm.org/D80450
2020-12-02 16:33:33 -05:00
Yaxun (Sam) Liu 5c8911d0ba [CUDA][HIP] Diagnose reference of host variable
This patch diagnoses invalid references of global host variables in device,
global, or host device functions.

Differential Revision: https://reviews.llvm.org/D91281
2020-12-02 10:15:56 -05:00
Artem Belevich be86b6773b [CUDA] Allow local static variables with target attributes.
While CUDA documentation claims that such variables are not allowed[1], NVCC has
been accepting them since CUDA-10.0[2] and some headers in CUDA-11 rely on this
working.

1. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#static-variables-function
2. https://godbolt.org/z/zsodzc

Differential Revision: https://reviews.llvm.org/D88345
2020-11-03 10:30:38 -08:00
Artem Belevich 0a3ebb4d8d Revert "[CUDA] Allow local static variables with target attributes."
This reverts commit f38a9e5117
Which triggered assertions.
2020-11-02 15:09:07 -08:00
Artem Belevich f38a9e5117 [CUDA] Allow local static variables with target attributes.
While CUDA documentation claims that such variables are not allowed[1], NVCC has
been accepting them since CUDA-10.0[2] and some headers in CUDA-11 rely on this
working.

1. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#static-variables-function
2. https://godbolt.org/z/zsodzc

Differential Revision: https://reviews.llvm.org/D88345
2020-11-02 14:37:13 -08:00
Yaxun (Sam) Liu 52bcd691cb Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"
This recommits 7f1f89ec8d and
40df06cdaf with bug fixes for
memory sanitizer failure and Tensile build failure.
2020-10-19 17:48:04 -04:00
Richard Smith f7f2e4261a PR47805: Use a single object for a function parameter in the caller and
callee in constant evaluation.

We previously made a deep copy of function parameters of class type when
passing them, resulting in the destructor for the parameter applying to
the original argument value, ignoring any modifications made in the
function body. This also meant that the 'this' pointer of the function
parameter could be observed changing between the caller and the callee.

This change completely reimplements how we model function parameters
during constant evaluation. We now model them roughly as if they were
variables living in the caller, albeit with an artificially reduced
scope that covers only the duration of the function call, instead of
modeling them as temporaries in the caller that we partially "reparent"
into the callee at the point of the call. This brings some minor
diagnostic improvements, as well as significantly reduced stack usage
during constant evaluation.
2020-10-14 17:43:51 -07:00
Richard Smith 69f7c006ff Revert "PR47805: Use a single object for a function parameter in the caller and"
Breaks a clangd unit test.

This reverts commit 8f8b9f2cca.
2020-10-13 19:32:03 -07:00
Richard Smith 8f8b9f2cca PR47805: Use a single object for a function parameter in the caller and
callee in constant evaluation.

We previously made a deep copy of function parameters of class type when
passing them, resulting in the destructor for the parameter applying to
the original argument value, ignoring any modifications made in the
function body. This also meant that the 'this' pointer of the function
parameter could be observed changing between the caller and the callee.

This change completely reimplements how we model function parameters
during constant evaluation. We now model them roughly as if they were
variables living in the caller, albeit with an artificially reduced
scope that covers only the duration of the function call, instead of
modeling them as temporaries in the caller that we partially "reparent"
into the callee at the point of the call. This brings some minor
diagnostic improvements, as well as significantly reduced stack usage
during constant evaluation.
2020-10-13 18:50:46 -07:00
Richard Smith ab870f3030 Revert "PR47805: Use a single object for a function parameter in the caller and"
The buildbots are displeased.

This reverts commit 8d03a972ce.
2020-10-13 15:59:00 -07:00
Richard Smith 8d03a972ce PR47805: Use a single object for a function parameter in the caller and
callee in constant evaluation.

We previously made a deep copy of function parameters of class type when
passing them, resulting in the destructor for the parameter applying to
the original argument value, ignoring any modifications made in the
function body. This also meant that the 'this' pointer of the function
parameter could be observed changing between the caller and the callee.

This change completely reimplements how we model function parameters
during constant evaluation. We now model them roughly as if they were
variables living in the caller, albeit with an artificially reduced
scope that covers only the duration of the function call, instead of
modeling them as temporaries in the caller that we partially "reparent"
into the callee at the point of the call. This brings some minor
diagnostic improvements, as well as significantly reduced stack usage
during constant evaluation.
2020-10-13 15:45:04 -07:00
Reid Kleckner 3453b6928d Revert "Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions""
This reverts commit e39da8ab6a.

This depends on a change that needs additional design review and needs
to be reverted.
2020-09-24 11:16:54 -07:00
Yaxun (Sam) Liu e39da8ab6a Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"
This recommits 7f1f89ec8d and
40df06cdaf after fixing memory
sanitizer failure.
2020-09-24 08:44:37 -04:00
Yaxun (Sam) Liu 772bd8a7d9 Revert "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"
This reverts commit 7f1f89ec8d.

This reverts commit 40df06cdaf.
2020-09-17 13:55:31 -04:00
Yaxun (Sam) Liu 40df06cdaf [CUDA][HIP] Defer overloading resolution diagnostics for host device functions
In CUDA/HIP a function may become implicit host device function by
pragma or constexpr. A host device function is checked in both
host and device compilation. However it may be emitted only
on host or device side, therefore the diagnostics should be
deferred until it is known to be emitted.

Currently clang is only able to defer certain diagnostics. This causes
false alarms and limits the usefulness of host device functions.

This patch lets clang defer all overloading resolution diagnostics for host device functions.

An option -fgpu-defer-diag is added to control this behavior. By default
it is off.

It is NFC for other languages.

Differential Revision: https://reviews.llvm.org/D84364
2020-09-17 11:30:42 -04:00
Yaxun (Sam) Liu 9275e14379 recommit 4fc752b30b [CUDA][HIP] Always defer diagnostics for wrong-sided reference
Fixed regression in test builtin-amdgcn-atomic-inc-dec-failure.cpp.
2020-07-17 09:14:39 -04:00
Yaxun (Sam) Liu a46ef7d42d Revert "[CUDA][HIP] Always defer diagnostics for wrong-sided reference"
This reverts commit 4fc752b30b.
2020-07-17 08:10:56 -04:00
Yaxun (Sam) Liu 4fc752b30b [CUDA][HIP] Always defer diagnostics for wrong-sided reference
When a device function calls a host function or vice versa, this is wrong-sided
reference. Currently clang immediately diagnose it. This is different from nvcc
behavior, where it is diagnosed only if the function is really emitted.

Current clang behavior causes false alarms for valid use cases.

This patch let clang always defer diagnostics for wrong-sided
reference.

Differential Revision: https://reviews.llvm.org/D83893
2020-07-17 07:51:55 -04:00
Yaxun (Sam) Liu 1eaad01046 [CUDA][HIP] Let lambda be host device by default
This patch let lambda be host device by default and adds diagnostics for
capturing host variable by reference in device lambda.

Differential Revision: https://reviews.llvm.org/D78655
2020-07-08 13:10:26 -04:00
Fangrui Song dfc0d94755 Revert D80450 "[CUDA][HIP] Fix implicit HD function resolution"
This reverts commit 263390d4f5.

This can still cause bogus errors:

eigen3/Eigen/src/Core/CoreEvaluators.h:94:38: error: call to implicitly-deleted copy constructor of 'unary_evaluator<Eigen::Inverse<Eigen::Matrix<double, 4, 4, 0, 4, 4>>>'

thrust/system/detail/generic/for_each.h:49:3: error: implicit instantiation of undefined template
'thrust::detail::STATIC_ASSERTION_FAILURE<false>'
2020-06-10 17:42:28 -07:00
Yaxun (Sam) Liu 263390d4f5 [CUDA][HIP] Fix implicit HD function resolution
recommit e03394c6a6 with fix

When implicit HD function calls a function in device compilation,
if one candidate is an implicit HD function, current resolution rule is:

D wins over HD and H
HD and H are equal

this caused regression when there is an otherwise worse D candidate

This patch changes that to

D, HD and H are all equal

The rationale is that we already know for host compilation there is already
a valid candidate in HD and H candidates that will not cause error. Allowing
HD and H gives us a fall back candidate that will not cause error. If D wins,
that means D has to be a better match otherwise, therefore D should also
be a valid candidate that will not cause error. In this way, we can guarantee
no regression.

Differential Revision: https://reviews.llvm.org/D80450
2020-06-04 16:54:52 -04:00
Yaxun (Sam) Liu 049d860707 [CUDA][HIP] Fix constexpr variables for C++17
constexpr variables are compile time constants and implicitly const, therefore
they are safe to emit on both device and host side. Besides, in many cases
they are intended for both device and host, therefore it makes sense
to emit them on both device and host sides if necessary.

In most cases constexpr variables are used as rvalue and the variables
themselves do not need to be emitted. However if their address is taken,
then they need to be emitted.

For C++14, clang is able to handle that since clang emits them with
available_externally linkage together with the initializer.

However for C++17, the constexpr static data member of a class or template class
become inline variables implicitly. Therefore they become definitions with
linkonce_odr or weak_odr linkages. As such, they can not have available_externally
linkage.

This patch fixes that by adding implicit constant attribute to
file scope constexpr variables and constexpr static data members
in device compilation.

Differential Revision: https://reviews.llvm.org/D79237
2020-06-03 21:56:52 -04:00
Artem Belevich ef649e8fd5 Revert "[CUDA][HIP] Workaround for resolving host device function against wrong-sided function"
Still breaks CUDA compilation.

This reverts commit e03394c6a6.
2020-05-18 12:22:55 -07:00
Yaxun (Sam) Liu e03394c6a6 [CUDA][HIP] Workaround for resolving host device function against wrong-sided function
recommit c77a4078e0 with fix

https://reviews.llvm.org/D77954 caused regressions due to diagnostics in implicit
host device functions.

For now, it seems the most feasible workaround is to treat implicit host device function and explicit host
device function differently. Basically in device compilation for implicit host device functions, keep the
old behavior, i.e. give host device candidates and wrong-sided candidates equal preference. For explicit
host device functions, favor host device candidates against wrong-sided candidates.

The rationale is that explicit host device functions are blessed by the user to be valid host device functions,
that is, they should not cause diagnostics in both host and device compilation. If diagnostics occur, user is
able to fix them. However, there is no guarantee that implicit host device function can be compiled in
device compilation, therefore we need to preserve its overloading resolution in device compilation.

Differential Revision: https://reviews.llvm.org/D79526
2020-05-12 08:27:50 -04:00
Artem Belevich bf6a26b066 Revert D77954 -- it breaks Eigen & Tensorflow.
This reverts commit 55bcb96f31.
2020-05-05 14:07:31 -07:00
Yaxun (Sam) Liu d75a6e93ae [CUDA][HIP] Fix empty ctor/dtor check for union
union ctor does not call ctors of its data members. union dtor does not call dtors of its data members.
Also union does not have base class.

Currently when clang checks whether union has an empty ctor/dtor, it checks the ctors/dtors of its
data members. This causes incorrectly diagnose device side global variables and shared variables as
having non-empty ctors/dtors.

This patch fixes that.

Differential Revision: https://reviews.llvm.org/D79367
2020-05-04 21:52:04 -04:00
Yaxun (Sam) Liu 55bcb96f31 recommit c77a4078e0 with fix
https://reviews.llvm.org/D77954 caused a regression about ambiguity of new operator
in file scope.

This patch recovered the previous behavior for comparison without a caller.

This is a workaround. For real fix we need D71227

https://reviews.llvm.org/D78970
2020-04-28 09:14:13 -04:00
Dmitri Gribenko 8c8aae852b Revert "recommit c77a4078e01033aa2206c31a579d217c8a07569b"
This reverts commit b46b1a916d. It broke
overload resolution for operator 'new' -- see reproducer in
https://reviews.llvm.org/D77954.
2020-04-27 16:41:35 +02:00
Yaxun (Sam) Liu b46b1a916d recommit c77a4078e0 2020-04-24 16:53:18 -04:00
Yaxun (Sam) Liu 7eae00477f Revert "[CUDA][HIP] Fix host/device based overload resolution"
This reverts commit c77a4078e0.
2020-04-24 14:57:10 -04:00