[clang] Extend diagnose_if to accept more detailed warning information, take 2 (#119712)

This is take two of #70976. This iteration of the patch makes sure that
custom
diagnostics without any warnin

[clang] Extend diagnose_if to accept more detailed warning information, take 2 (#119712)

This is take two of #70976. This iteration of the patch makes sure that
custom
diagnostics without any warning group don't get promoted by `-Werror` or
`-Wfatal-errors`.

This implements parts of the extension proposed in
https://discourse.llvm.org/t/exposing-the-diagnostic-engine-to-c/73092/7.

Specifically, this makes it possible to specify a diagnostic group in an
optional third argument.

show more ...

[Sema] Remove unused includes (NFC) (#116461)

Identified with misc-include-cleaner.

[clang] [NFC] Merge two ifs to a single one (#116226)

Revert "[clang] Extend diagnose_if to accept more detailed warning information (#70976)"

This reverts commit e39205654dc11c50bd117e8ccac243a641ebd71f.

There are further discussions in
https://githu

Revert "[clang] Extend diagnose_if to accept more detailed warning information (#70976)"

This reverts commit e39205654dc11c50bd117e8ccac243a641ebd71f.

There are further discussions in
https://github.com/llvm/llvm-project/pull/70976, happening for past two
weeks. Since there were no responses for couple weeks now, reverting
until author is back.

show more ...

Reapply "Reapply "[clang] Extend diagnose_if to accept more detailed warning information (#70976)" (#108453)"

This reverts commit e1bd9740faa62c11cc785a7b70ec1ad17e286bd1.

Fixes incorrect use of th

Reapply "Reapply "[clang] Extend diagnose_if to accept more detailed warning information (#70976)" (#108453)"

This reverts commit e1bd9740faa62c11cc785a7b70ec1ad17e286bd1.

Fixes incorrect use of the `DiagnosticsEngine` in the clangd tests.

show more ...

Revert "Reapply "[clang] Extend diagnose_if to accept more detailed warning information (#70976)" (#108453)"

This reverts commit e7f782e7481cea23ef452a75607d3d61f5bd0d22.

This had UBSan failures:

Revert "Reapply "[clang] Extend diagnose_if to accept more detailed warning information (#70976)" (#108453)"

This reverts commit e7f782e7481cea23ef452a75607d3d61f5bd0d22.

This had UBSan failures:

[----------] 1 test from ConfigCompileTests
[ RUN ] ConfigCompileTests.DiagnosticSuppression
Config fragment: compiling <unknown>:0 -> 0x00007B8366E2F7D8 (trusted=false)
/usr/local/google/home/fmayer/large/llvm-project/llvm/include/llvm/ADT/IntrusiveRefCntPtr.h:203:33: runtime error: reference binding to null pointer of type 'clang::DiagnosticIDs'

UndefinedBehaviorSanitizer: undefined-behavior /usr/local/google/home/fmayer/large/llvm-project/llvm/include/llvm/ADT/IntrusiveRefCntPtr.h:203:33

Pull Request: https://github.com/llvm/llvm-project/pull/108645

show more ...

Reapply "[clang] Extend diagnose_if to accept more detailed warning information (#70976)" (#108453)

This reverts commit e0cd11eba526234ca14a0b91f5598ca3363b6aca.

Update the use of `getWarningOpti

Reapply "[clang] Extend diagnose_if to accept more detailed warning information (#70976)" (#108453)

This reverts commit e0cd11eba526234ca14a0b91f5598ca3363b6aca.

Update the use of `getWarningOptionForDiag` in flang to use the
DiagnosticIDs.

show more ...

Revert "[clang] Extend diagnose_if to accept more detailed warning information (#70976)"

This reverts commit 030c6da7af826b641db005be925b20f956c3a6bb.

Several build bots are failing:
https://lab.ll

Revert "[clang] Extend diagnose_if to accept more detailed warning information (#70976)"

This reverts commit 030c6da7af826b641db005be925b20f956c3a6bb.

Several build bots are failing:
https://lab.llvm.org/buildbot/#/builders/89/builds/6211
https://lab.llvm.org/buildbot/#/builders/157/builds/7578
https://lab.llvm.org/buildbot/#/builders/140/builds/6429

show more ...

[clang] Extend diagnose_if to accept more detailed warning information (#70976)

This implements parts of the extension proposed in
https://discourse.llvm.org/t/exposing-the-diagnostic-engine-to-c/7

[clang] Extend diagnose_if to accept more detailed warning information (#70976)

This implements parts of the extension proposed in
https://discourse.llvm.org/t/exposing-the-diagnostic-engine-to-c/73092/7.

Specifically, this makes it possible to specify a diagnostic group in an
optional third argument.

show more ...

[Offload][CUDA] Allow CUDA kernels to use LLVM/Offload (#94549)

Through the new `-foffload-via-llvm` flag, CUDA kernels can now be
lowered to the LLVM/Offload API. On the Clang side, this is simply

[Offload][CUDA] Allow CUDA kernels to use LLVM/Offload (#94549)

Through the new `-foffload-via-llvm` flag, CUDA kernels can now be
lowered to the LLVM/Offload API. On the Clang side, this is simply done
by using the OpenMP offload toolchain and emitting calls to `llvm*`
functions to orchestrate the kernel launch rather than `cuda*`
functions. These `llvm*` functions are implemented on top of the
existing LLVM/Offload API.

As we are about to redefine the Offload API, this wil help us in the
design process as a second offload language.

We do not support any CUDA APIs yet, however, we could:
https://www.osti.gov/servlets/purl/1892137

For proper host execution we need to resurrect/rebase
https://tianshilei.me/wp-content/uploads/2021/12/llpp-2021.pdf
(which was designed for debugging).

```
❯❯❯ cat test.cu
extern "C" {
void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
}

__global__ void square(int *A) { *A = 42; }

int main(int argc, char **argv) {
int DevNo = 0;
int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
*Ptr = 7;
printf("Ptr %p, *Ptr %i\n", Ptr, *Ptr);
square<<<1, 1>>>(Ptr);
printf("Ptr %p, *Ptr %i\n", Ptr, *Ptr);
llvm_omp_target_free_shared(Ptr, DevNo);
}

❯❯❯ clang++ test.cu -O3 -o test123 -foffload-via-llvm --offload-arch=native

❯❯❯ llvm-objdump --offloading test123

test123: file format elf64-x86-64

OFFLOADING IMAGE [0]:
kind elf
arch gfx90a
triple amdgcn-amd-amdhsa
producer openmp

❯❯❯ LIBOMPTARGET_INFO=16 ./test123
Ptr 0x155448ac8000, *Ptr 7
Ptr 0x155448ac8000, *Ptr 42
```

show more ...

[CUDA][HIP] warn incompatible redeclare (#77359)

nvcc warns about the following code:

`void f();
__device__ void f() {}`

but clang does not since clang allows device function to overload host

[CUDA][HIP] warn incompatible redeclare (#77359)

nvcc warns about the following code:

`void f();
__device__ void f() {}`

but clang does not since clang allows device function to overload host
function.

Users want clang to emit similar warning to help code to be compatible
with nvcc.

Since this may cause regression with existing code, the warning is off
by default and can be enabled by -Wnvcc-compat.

It won't cause warning in system headers, even with -Wnvcc-compat.

show more ...

[clang] Introduce `SemaCUDA` (#88559)

This patch moves CUDA-related `Sema` function into new `SemaCUDA` class,
following the recent example of SYCL, OpenACC, and HLSL. This is a part
of the effort

[clang] Introduce `SemaCUDA` (#88559)

This patch moves CUDA-related `Sema` function into new `SemaCUDA` class,
following the recent example of SYCL, OpenACC, and HLSL. This is a part
of the effort to split Sema. Additional context can be found in
https://github.com/llvm/llvm-project/pull/82217,
https://github.com/llvm/llvm-project/pull/84184,
https://github.com/llvm/llvm-project/pull/87634.

show more ...

[clang][NFC] Refactor `CUDAFunctionTarget`

Refactor `CUDAFunctionTarget` into a scoped enum at namespace scope, so that it can be forward declared. This is done in preparation for `SemaCUDA`.

[clang][NFC] Refactor `CXXSpecialMember`

In preparation for `SemaCUDA`, which requires this enum to be forward-declarable.

[HIP] fix host-used external kernel (#83870)

In -fgpu-rdc mode, when an external kernel is used by a host function
with weak_odr linkage (e.g. explicitly instantiated template function),
the kerne

[HIP] fix host-used external kernel (#83870)

In -fgpu-rdc mode, when an external kernel is used by a host function
with weak_odr linkage (e.g. explicitly instantiated template function),
the kernel should not be marked as host-used external kernel, since the
host function may be dropped by the linker. Mark the external kernel as
host-used external kernel will force a reference to the external kernel,
which the user may not define in other TU.

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

show more ...

[CUDA][HIP] allow trivial ctor/dtor in device var init (#73140)

Treat ctor/dtor in device var init as host device function
so that they can be used to initialize file-scope
device variables to mat

[CUDA][HIP] allow trivial ctor/dtor in device var init (#73140)

Treat ctor/dtor in device var init as host device function
so that they can be used to initialize file-scope
device variables to match nvcc behavior. If they are non-trivial
they will be diagnosed.

We cannot add implicit host device attrs to non-trivial
ctor/dtor since determining whether they are non-trivial
needs to know whether they have a trivial body and all their
member and base classes' ctor/dtor have trivial body, which
is affected by where their bodies are defined or instantiated.

Fixes: #72261

Fixes: SWDEV-432412

show more ...

Revert "[CUDA][HIP] make trivial ctor/dtor host device (#72394)"

This reverts commit 27e6e4a4d0e3296cebad8db577ec0469a286795e.

This patch is reverted due to regression. A testcase is:

`template <c

Revert "[CUDA][HIP] make trivial ctor/dtor host device (#72394)"

This reverts commit 27e6e4a4d0e3296cebad8db577ec0469a286795e.

This patch is reverted due to regression. A testcase is:

`template <class T>
struct ptr {
~ptr() { static int x = 1;}
};

template <class T>
struct Abc : ptr<T> {
public:
Abc();
~Abc() {}
};

template
class Abc<int>;
`

show more ...

[CUDA][HIP] make trivial ctor/dtor host device (#72394)

Make trivial ctor/dtor implicitly host device functions so that they can
be used to initialize file-scope
device variables to match nvcc beh

[CUDA][HIP] make trivial ctor/dtor host device (#72394)

Make trivial ctor/dtor implicitly host device functions so that they can
be used to initialize file-scope
device variables to match nvcc behavior.

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

Fixes: SWDEV-432412

show more ...

[CUDA][HIP] Make template implicitly host device (#70369)

Added option -foffload-implicit-host-device-templates which is off by
default.

When the option is on, template functions and specializat

[CUDA][HIP] Make template implicitly host device (#70369)

Added option -foffload-implicit-host-device-templates which is off by
default.

When the option is on, template functions and specializations without
host/device attributes have implicit host device attributes.

They can be overridden by device template functions with the same
signagure.
They are emitted on device side only if they are used on device side.

This feature is added as an extension.
`__has_extension(cuda_implicit_host_device_templates)` can be used to
check whether it is enabled.

This is to facilitate using standard C++ headers for device.

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

Fixes: SWDEV-428314

show more ...

[CUDA][HIP] Fix init var diag in temmplate (#69081)

Currently clang diagnoses the following code:
(https://godbolt.org/z/s8zK3E5P5) but nvcc
does not.

`
struct A {
constexpr A(){}
};

s

[CUDA][HIP] Fix init var diag in temmplate (#69081)

Currently clang diagnoses the following code:
(https://godbolt.org/z/s8zK3E5P5) but nvcc
does not.

`
struct A {
constexpr A(){}
};

struct B {
A a;
int b;
};

template<typename T>
__global__ void kernel( )
{
__shared__ B x;
}
`

Clang generates an implicit trivial ctor for struct B, which should be
allowed for initializing a shared variable.

However, the body of the ctor is defined only if the template kernel is
instantiated. Clang checks the initialization of variable in
non-instantiated templates, where it cannot find the body of the ctor,
therefore diagnoses it.

This patch skips the check for non-instantiated templates.

show more ...

[Clang] Handle consteval expression in array bounds expressions (#66222)

The bounds of a c++ array is a _constant-expression_. And in C++ it is
also a constant expression.

But we also support VL

[Clang] Handle consteval expression in array bounds expressions (#66222)

The bounds of a c++ array is a _constant-expression_. And in C++ it is
also a constant expression.

But we also support VLAs, ie arrays with non-constant bounds.

We need to take care to handle the case of a consteval function (which
are specified to be only immediately called in non-constant contexts)
that appear in arrays bounds.

This introduces `Sema::isAlwayConstantEvaluatedContext`, and a flag in
ExpressionEvaluationContextRecord, such that immediate functions in
array bounds are always immediately invoked.

Sema had both `isConstantEvaluatedContext` and
`isConstantEvaluated`, so I took the opportunity to cleanup that.

The change in `TimeProfilerTest.cpp` is an unfortunate manifestation of
the problem that #66203 seeks to address.

Fixes #65520

show more ...

[HIP][Clang][Sema] Add Sema support for `hipstdpar`

This patch adds the Sema changes needed for enabling HIP parallel algorithm offload on AMDGPU targets. This change impacts the CUDA / HIP language

[HIP][Clang][Sema] Add Sema support for `hipstdpar`

This patch adds the Sema changes needed for enabling HIP parallel algorithm offload on AMDGPU targets. This change impacts the CUDA / HIP language specific checks, and only manifests if compiling in `hipstdpar` mode. In this case, we essentially do three things:

1. Allow device side callers to call host side callees - since the user visible HLL would be standard C++, with no annotations / restriction mechanisms, we cannot unambiguously establish that such a call is an error, so we conservatively allow all such calls, deferring actual cleanup to a subsequent pass over IR;
2. Allow host formed lambdas to capture by reference;
3. Allow device functions to use host global variables.

Reviewed by: yaxunl

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

show more ...

Reland "[CUDA][HIP] Fix overloading resolution in global var init" (#65606)

https://reviews.llvm.org/D158247 caused regressions for HIP on Windows
and was reverted.

A reduced test case is:

``

Reland "[CUDA][HIP] Fix overloading resolution in global var init" (#65606)

https://reviews.llvm.org/D158247 caused regressions for HIP on Windows
and was reverted.

A reduced test case is:

```
typedef void (__stdcall* funcTy)();
void invoke(funcTy f);

static void __stdcall callee() noexcept {
}

void foo() {
invoke(callee);
}
```

It is due to clang missing handling host/device attributes for calling
convention at a few places

This patch fixes that.

show more ...

Revert "[CUDA][HIP] Fix overloading resolution in global variable initializer"

This reverts commit de0df639724b10001ea9a74539381ea494296be9.

It was reverted due to regression in HIP unit test on Wi

Revert "[CUDA][HIP] Fix overloading resolution in global variable initializer"

This reverts commit de0df639724b10001ea9a74539381ea494296be9.

It was reverted due to regression in HIP unit test on Windows:

In file included from C:\hip-tests\catch\unit\graph\hipGraphClone.cc:37:

In file included from C:\hip-tests\catch\.\include\hip_test_common.hh:24:

In file included from C:\hip-tests\catch\.\include/hip_test_context.hh:24:

In file included from C:/install/native/Release/x64/hip/include\hip/hip_runtime.h:54:

C:/dk/win\vc\14.31.31107\include\thread:76:70: error: cannot initialize a parameter of type '_beginthreadex_proc_type' (aka 'unsigned int (*)(void *) __attribute__((stdcall))') with an lvalue of type 'const unsigned int (*)(void *) noexcept __attribute__((stdcall))': different exception specifications

76 | reinterpret_cast<void*>(_CSTD _beginthreadex(nullptr, 0, _Invoker_proc, _Decay_copied.get(), 0, &_Thr._Id));

| ^~~~~~~~~~~~~

C:\hip-tests\catch\unit\graph\hipGraphClone.cc:290:21) &>' requested here

90 | _Start(_STD forward<_Fn>(_Fx), _STD forward<_Args>(_Ax)...);

| ^

C:\hip-tests\catch\unit\graph\hipGraphClone.cc:290:21) &, 0>' requested here

311 | std::thread t(lambdaFunc);

| ^

C:/dk/win\ms_wdk\e22621\Include\10.0.22621.0\ucrt\process.h:99:40: note: passing argument to parameter '_StartAddress' here

99 | _In_ _beginthreadex_proc_type _StartAddress,

| ^

1 error generated when compiling for gfx1030.

show more ...

[CUDA][HIP] Fix overloading resolution in global variable initializer

Currently, clang does not resolve certain overloaded functions correctly in the initializer
of global variables, e.g.

template<

[CUDA][HIP] Fix overloading resolution in global variable initializer

Currently, clang does not resolve certain overloaded functions correctly in the initializer
of global variables, e.g.

template<typename T1, typename U>
T1 mypow(T1, U);

__attribute__((device)) double mypow(double, int);

double t_extent = mypow(1.0, 2);

In the above example, mypow is supposed to resolve to the host version
but clang resolves it to the device version instead, and emits an error
(https://godbolt.org/z/17xxzaa67).

However, if the variable is assigned in a host function, there is no error.
The discrepancy in overloading resolution inside and outside of
a function is due to clang not accounting for the host/device target
when resolving functions called in the initializer of a global variable.

This patch introduces a global host/device target context for CUDA/HIP
for functions called outside of functions. For global variable initialization,
it is determined by the host/device attribute of the variable. For other
situations, a default value of host_device is sufficient.

Reviewed by: Artem Belevich

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

Fixes: SWDEV-416731

show more ...