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

[Artem Belevich via cfe-commits]

================
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:   -isystem %S/Inputs -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:   -isystem %S/Inputs -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:   -isystem %S/Inputs -verify=redecl -Woffload-incompatible-redeclare %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:   -isystem %S/Inputs -fcuda-is-device -Woffload-incompatible-redeclare -verify=redecl %s
+
+// expected-no-diagnostics
+#include "cuda.h"
+
+__device__ void f(); // redecl-note {{previous declaration is here}}
+
+void f() {} // redecl-warning {{incompatible host/device attribute with redeclaration: new declaration is __host__ function, old declaration is __device__ function. It will cause warning with nvcc}}
----------------
Artem-B wrote:

> "this is incompatible with X" warning like we have for GCC

CUDA compilation includes a lot of CUDA SDK and host headers. `<cmath>` is one of them. Compiler will *always* see a lot of implicitly `__host__` math functions declared/defined by the host headers and a lot of the same functions with `__device__` attribute declared/defined for the GPU. Whether or not user code has the same warning becomes rather moot, when *all* compilations will produce tons of these warnings. 

That said, we *may* be saved by the fact that (IIRC) compiler hides the warnings that come from the system headers.

It would be good to test it on a real CUDA compilation first. If we're guaranteed to have dozens of those warnings for every compilation, it would render the flag practically useless.


https://github.com/llvm/llvm-project/pull/77359