[CUDA] remove duplicate condition

Reviewed by: Yaxun Liu

Differential Revision: https://reviews.llvm.org/D130168

Change-Id: Ia00c3dfa9ea20e61235817fd4bb61d33c7c98a60

[clang] Use value instead of getValue (NFC)

[clang] Don't use Optional::hasValue (NFC)

This patch replaces Optional::hasValue with the implicit cast to bool
in conditionals only.

Revert "Don't use Optional::hasValue (NFC)"

This reverts commit aa8feeefd3ac6c78ee8f67bf033976fc7d68bc6d.

Don't use Optional::hasValue (NFC)

[clang] Don't use Optional::hasValue (NFC)

[CUDA][HIP] Fix host used external kernel in archive

For -fgpu-rdc, a host function may call an external kernel
which is defined in an archive of bitcode. Since this external
kernel is only referenc

[CUDA][HIP] Fix host used external kernel in archive

For -fgpu-rdc, a host function may call an external kernel
which is defined in an archive of bitcode. Since this external
kernel is only referenced in host function, the device
bitcode does not contain reference to this external
kernel, then the linker will not try to resolve this external
kernel in the archive.

To fix this issue, host-used external kernels and device
variables are tracked. A global array containing pointers
to these external kernels and variables is emitted which
serves as an artificial references to the external kernels
and variables used by host.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D123441

show more ...

[CUDA][HIP] Fix hostness check with -fopenmp

CUDA/HIP determines whether a function can be called based on
the device/host attributes of callee and caller. Clang assumes the
caller is CurContext. Th

[CUDA][HIP] Fix hostness check with -fopenmp

CUDA/HIP determines whether a function can be called based on
the device/host attributes of callee and caller. Clang assumes the
caller is CurContext. This is correct in most cases, however, it is
not correct in OpenMP parallel region when CUDA/HIP program
is compiled with -fopenmp. This causes incorrect overloading
resolution and missed diagnostics.

To get the correct caller, clang needs to chase the parent chain
of DeclContext starting from CurContext until a function decl
or a lambda decl is reached. Sema API is adapted to achieve that
and used to determine the caller in hostness check.

Reviewed by: Artem Belevich, Richard Smith

Differential Revision: https://reviews.llvm.org/D121765

show more ...

Use llvm::append_range instead of push_back loops where applicable. NFCI.

[CUDA][HIP] Do not promote constexpr var with non-constant initializer

constexpr var may be initialized with address of non-const variable.
In this case the initializer is not constant in device com

[CUDA][HIP] Do not promote constexpr var with non-constant initializer

constexpr var may be initialized with address of non-const variable.
In this case the initializer is not constant in device compilation.
This has been handled for const vars but not for constexpr vars.

This patch makes handling of const var and constexpr var
consistent.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D119615

Fixes: https://github.com/llvm/llvm-project/issues/53780

show more ...

[CUDA][HIP] Do not treat host var address as constant in device compilation

Currently clang treats host var address as constant in device compilation,
which causes const vars initialized with host v

[CUDA][HIP] Do not treat host var address as constant in device compilation

Currently clang treats host var address as constant in device compilation,
which causes const vars initialized with host var address promoted to
device variables incorrectly and results in undefined symbols.

This patch fixes that.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D118153

Fixes: SWDEV-309881

Change-Id: I0a69357063c6f8539ef259c96c250d04615f4473

show more ...

Remove redundant return and continue statements (NFC)

Identified with readability-redundant-control-flow.

[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

[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

show more ...

[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 e

[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

show more ...

[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 f

[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

show more ...

[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

D

[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

show more ...

[CUDA] Normalize handling of defauled dtor.

Defaulted destructor was treated inconsistently, compared to other
compiler-generated functions.

When Sema::IdentifyCUDATarget() got called on just-creat

[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

show more ...

[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.

[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

show more ...

[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 i

[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

show more ...

Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

This recommits 7f1f89ec8d9944559042bb6d3b1132eabe3409de and
40df06cdafc010002fc9cfe1dda73d689b7d27a6 with bu

Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

This recommits 7f1f89ec8d9944559042bb6d3b1132eabe3409de and
40df06cdafc010002fc9cfe1dda73d689b7d27a6 with bug fixes for
memory sanitizer failure and Tensile build failure.

show more ...

Revert "Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions""

This reverts commit e39da8ab6a286ac777d5fe7799f1eb782cf99938.

This depends on a change that needs

Revert "Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions""

This reverts commit e39da8ab6a286ac777d5fe7799f1eb782cf99938.

This depends on a change that needs additional design review and needs
to be reverted.

show more ...

Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

This recommits 7f1f89ec8d9944559042bb6d3b1132eabe3409de and
40df06cdafc010002fc9cfe1dda73d689b7d27a6 after f

Recommit "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

This recommits 7f1f89ec8d9944559042bb6d3b1132eabe3409de and
40df06cdafc010002fc9cfe1dda73d689b7d27a6 after fixing memory
sanitizer failure.

show more ...

Revert "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

This reverts commit 7f1f89ec8d9944559042bb6d3b1132eabe3409de.

This reverts commit 40df06cdafc010002fc9cfe1dda

Revert "[CUDA][HIP] Defer overloading resolution diagnostics for host device functions"

This reverts commit 7f1f89ec8d9944559042bb6d3b1132eabe3409de.

This reverts commit 40df06cdafc010002fc9cfe1dda73d689b7d27a6.

show more ...

[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 check

[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

show more ...

recommit 4fc752b30b9a [CUDA][HIP] Always defer diagnostics for wrong-sided reference

Fixed regression in test builtin-amdgcn-atomic-inc-dec-failure.cpp.