[CUDA][HIP] warn incompatible redeclare #77359

yxsamliu · 2024-01-08T19:08:02Z

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.

llvmbot · 2024-01-08T19:08:31Z

@llvm/pr-subscribers-clang

Author: Yaxun (Sam) Liu (yxsamliu)

Changes

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 -Woffload-incompatible-redeclare.

It won't cause warning in system headers, even with -Woffload-incompatible-redeclare.

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

3 Files Affected:

(modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+6)
(modified) clang/lib/Sema/SemaCUDA.cpp (+21-14)
(added) clang/test/SemaCUDA/function-redclare.cu (+19)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index a97182cad5d513..a4c5a76b59cb1b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8837,6 +8837,12 @@ def err_cuda_ovl_target : Error<
   "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;
 def note_cuda_ovl_candidate_target_mismatch : Note<
     "candidate template ignored: target attributes do not match">;
+def warn_offload_incompatible_redeclare : Warning<
+  "incompatible host/device attribute with redeclaration: "
+  "new declaration is %select{__device__|__global__|__host__|__host__ __device__}0 function, "
+  "old declaration is %select{__device__|__global__|__host__|__host__ __device__}1 function. "
+  "It will cause warning with nvcc">,
+  InGroup<DiagGroup<"offload-incompatible-redeclare">>, DefaultIgnore;
 
 def err_cuda_device_builtin_surftex_cls_template : Error<
     "illegal device builtin %select{surface|texture}0 reference "
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 6a66ecf6f94c17..27da1775d4751d 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -992,22 +992,29 @@ void Sema::checkCUDATargetOverload(FunctionDecl *NewFD,
     // HD/global functions "exist" in some sense on both the host and device, so
     // should have the same implementation on both sides.
     if (NewTarget != OldTarget &&
-        ((NewTarget == CFT_HostDevice &&
-          !(LangOpts.OffloadImplicitHostDeviceTemplates &&
-            isCUDAImplicitHostDeviceFunction(NewFD) &&
-            OldTarget == CFT_Device)) ||
-         (OldTarget == CFT_HostDevice &&
-          !(LangOpts.OffloadImplicitHostDeviceTemplates &&
-            isCUDAImplicitHostDeviceFunction(OldFD) &&
-            NewTarget == CFT_Device)) ||
-         (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) &&
         !IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
                     /* ConsiderCudaAttrs = */ false)) {
-      Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
-          << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
-      Diag(OldFD->getLocation(), diag::note_previous_declaration);
-      NewFD->setInvalidDecl();
-      break;
+      if ((NewTarget == CFT_HostDevice &&
+           !(LangOpts.OffloadImplicitHostDeviceTemplates &&
+             isCUDAImplicitHostDeviceFunction(NewFD) &&
+             OldTarget == CFT_Device)) ||
+          (OldTarget == CFT_HostDevice &&
+           !(LangOpts.OffloadImplicitHostDeviceTemplates &&
+             isCUDAImplicitHostDeviceFunction(OldFD) &&
+             NewTarget == CFT_Device)) ||
+          (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) {
+        Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
+            << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
+        Diag(OldFD->getLocation(), diag::note_previous_declaration);
+        NewFD->setInvalidDecl();
+        break;
+      }
+      if ((NewTarget == CFT_Host && OldTarget == CFT_Device) ||
+          (NewTarget == CFT_Device && OldTarget == CFT_Host)) {
+        Diag(NewFD->getLocation(), diag::warn_offload_incompatible_redeclare)
+            << NewTarget << OldTarget;
+        Diag(OldFD->getLocation(), diag::note_previous_declaration);
+      }
     }
   }
 }
diff --git a/clang/test/SemaCUDA/function-redclare.cu b/clang/test/SemaCUDA/function-redclare.cu
new file mode 100644
index 00000000000000..266483ebf49f0e
--- /dev/null
+++ b/clang/test/SemaCUDA/function-redclare.cu
@@ -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}}
+
+void g(); // redecl-note {{previous declaration is here}}
+
+__device__ void g() {} // redecl-warning {{incompatible host/device attribute with redeclaration: new declaration is __device__ function, old declaration is __host__ function. It will cause warning with nvcc}}

yxsamliu · 2024-01-16T15:14:22Z

ping

Artem-B · 2024-01-16T19:33:18Z

clang/test/SemaCUDA/function-redclare.cu

+__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}}


I'm not convinced that it's something we need to fix. As far as CUDA is concerned those are two different function overloads and it's business as usual with no warning required.

nvcc does not have a concept of attribute-based overloads, so for them a function with different attributes is a sign of a potential issue, but it's not the case for clang, IMO.

I believe turning the warning on will never work in practice, because there will likely almost always be a ton of math functions with host and device overloads.
As I said, it is very common to have those and a lot of such overloads get pulled in from the headers pre-included by clang.

IMO that looks like a nice "this is incompatible with X" warning like we have for GCC, different C/C++ versions etc., but I'm not an expert here, so maybe this isn't actually an incompatibility? FWIW neither the diagnostic nor your comment make this really clear to me. (Also I'd drop the It will cause warning with nvcc and make the flag something like -Wnvcc-compat)

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

OK. If there's a precedent, it makes things easier.

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.

I checked with real HIP apps and the warnings will only show up for user's code or header files. There are no warnings for host/device redeclarations in HIP or clang headers since they are included as system headers.

IMO that looks like a nice "this is incompatible with X" warning like we have for GCC, different C/C++ versions etc., but I'm not an expert here, so maybe this isn't actually an incompatibility? FWIW neither the diagnostic nor your comment make this really clear to me. (Also I'd drop the It will cause warning with nvcc and make the flag something like -Wnvcc-compat)

sorry. I missed your comments. will rename the flag

yxsamliu · 2024-06-06T13:56:18Z

ping

Our users keep requesting this feature since they want their HIP code works with both nvcc and clang. I tested it with real HIP apps and did not see warnings emitted for clang wrapper headers and HIP system headers. Only warnings for users' own code were emitted. Also since this warning is off by default, it won't affect normal users.

Artem-B

LGTM with some wording/namiung nits.

Artem-B · 2024-06-06T20:51:08Z

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9013,6 +9013,12 @@ def err_cuda_ovl_target : Error<
  "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;
 def note_cuda_ovl_candidate_target_mismatch : Note<
    "candidate template ignored: target attributes do not match">;
+def warn_offload_incompatible_redeclare : Warning<


We should have some sort of umbrella warning option for nvcc compatibility. Function overloads are the primary source of the differences, but we have other differences that the users may want to know about. E.g. some of the compiler builtins would be different. There are probably other things.

This option is fine, for now.

Also, we should document it.

will document it

Artem-B · 2024-06-06T20:55:26Z

clang/include/clang/Basic/DiagnosticSemaKinds.td

+  "incompatible host/device attribute with redeclaration: "
+  "new declaration is %select{__device__|__global__|__host__|__host__ __device__}0 function, "
+  "old declaration is %select{__device__|__global__|__host__|__host__ __device__}1 function. "
+  "It will cause warning with nvcc">,


It will potentially be a more serious issue, than a mere warning. the source code has both functions, it's possible that the implementation lives in different TUs. With clang, they will be treated as function overloads and things will work. With NVCC, they will end up being treated as GPU functions and that will result in potential ORD violation becaue the user will end up with two different GPU-side functions. It will be a problem during RDC compilation which may link both instances into a single GPU executable.

I'd rephrase it in more general terms. Maybe something along the lines of "Target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration". We should tell what we're diagnosing, but make no opinion on whether it will be a problem in any specific case.

will modify the diagnostic message

nvcc warns about the following code: 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 -Woffload-incompatible-redeclare. It won't cause warning in system headers, even with -Woffload-incompatible-redeclare.

reverts: composableKernels 53d2f4d [CUDA][HIP] warn incompatible redeclare (llvm#77359) breaks MIOpen 18ec885 [RFC][AMDGPU] Remove old llvm.amdgcn.buffer.* and tbuffer intrinsics (llvm#93801) Change-Id: I7191261fb80fe1dc8a47063f74d785ab415290b0

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.

This PR was already merged in trunk: llvm#77359 However, it was reverted in amd-staging due to ck issue ROCm/composable_kernel#1330 Now try relanding it in amd-staging after ck issue fixed by ROCm/composable_kernel#1342 Fixes: SWDEV-431838 Original commit message: 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. Change-Id: Ia370700eb3eb1b1928d04ec59e2ec63506f85545

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.

yxsamliu requested a review from Artem-B January 8, 2024 19:08

llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Jan 8, 2024

Artem-B reviewed Jan 16, 2024

View reviewed changes

yxsamliu force-pushed the warn-redecl branch from 30daa34 to 319552f Compare May 30, 2024 20:05

Artem-B approved these changes Jun 6, 2024

View reviewed changes

yxsamliu force-pushed the warn-redecl branch from 319552f to dd58965 Compare June 7, 2024 16:59

rename warning

d81c53c

yxsamliu merged commit 53d2f4d into llvm:main Jun 10, 2024
8 checks passed

yxsamliu mentioned this pull request Jun 11, 2024

need to add -Wno-nvcc-compt ROCm/composable_kernel#1330

Closed

HerrCai0907 mentioned this pull request Jun 13, 2024

tidy #95384

Closed

searlmc1 pushed a commit to ROCm/llvm-project that referenced this pull request Apr 15, 2025

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

3450a62

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[CUDA][HIP] warn incompatible redeclare #77359

[CUDA][HIP] warn incompatible redeclare #77359

Uh oh!

yxsamliu commented Jan 8, 2024 •

edited

Loading

Uh oh!

llvmbot commented Jan 8, 2024

Uh oh!

yxsamliu commented Jan 16, 2024

Uh oh!

Artem-B Jan 16, 2024

Uh oh!

Artem-B Jan 16, 2024

Uh oh!

philnik777 Jan 16, 2024 •

edited

Loading

Uh oh!

Artem-B Jan 16, 2024 •

edited

Loading

Uh oh!

yxsamliu May 31, 2024

Uh oh!

yxsamliu Jun 7, 2024

Uh oh!

yxsamliu commented Jun 6, 2024

Uh oh!

Artem-B left a comment

Uh oh!

Artem-B Jun 6, 2024

Uh oh!

yxsamliu Jun 7, 2024

Uh oh!

Artem-B Jun 6, 2024

Uh oh!

yxsamliu Jun 7, 2024

Uh oh!

Uh oh!

Uh oh!

		__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}}

[CUDA][HIP] warn incompatible redeclare #77359

[CUDA][HIP] warn incompatible redeclare #77359

Uh oh!

Conversation

yxsamliu commented Jan 8, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jan 8, 2024

Uh oh!

yxsamliu commented Jan 16, 2024

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

philnik777 Jan 16, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Artem-B Jan 16, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yxsamliu commented Jun 6, 2024

Uh oh!

Artem-B left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

yxsamliu commented Jan 8, 2024 •

edited

Loading

philnik777 Jan 16, 2024 •

edited

Loading

Artem-B Jan 16, 2024 •

edited

Loading