Skip to content

[CUDA][HIP] warn incompatible redeclare #77359

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Jun 10, 2024
Merged

[CUDA][HIP] warn incompatible redeclare #77359

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
dd58965
[CUDA][HIP] warn incompatible redeclare
yxsamliu May 30, 2024
d81c53c
rename warning
yxsamliu Jun 7, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9013,6 +9013,11 @@ 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<
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

will document it

"target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:"
"new declaration is %select{__device__|__global__|__host__|__host__ __device__}0 function, "
"old declaration is %select{__device__|__global__|__host__|__host__ __device__}1 function">,
InGroup<DiagGroup<"nvcc-compat">>, DefaultIgnore;

def err_cuda_device_builtin_surftex_cls_template : Error<
"illegal device builtin %select{surface|texture}0 reference "
Expand Down
41 changes: 25 additions & 16 deletions clang/lib/Sema/SemaCUDA.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1018,24 +1018,33 @@ void SemaCUDA::checkTargetOverload(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 == CUDAFunctionTarget::HostDevice &&
!(getLangOpts().OffloadImplicitHostDeviceTemplates &&
isImplicitHostDeviceFunction(NewFD) &&
OldTarget == CUDAFunctionTarget::Device)) ||
(OldTarget == CUDAFunctionTarget::HostDevice &&
!(getLangOpts().OffloadImplicitHostDeviceTemplates &&
isImplicitHostDeviceFunction(OldFD) &&
NewTarget == CUDAFunctionTarget::Device)) ||
(NewTarget == CUDAFunctionTarget::Global) ||
(OldTarget == CUDAFunctionTarget::Global)) &&
!SemaRef.IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
/* ConsiderCudaAttrs = */ false)) {
Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
<< llvm::to_underlying(NewTarget) << NewFD->getDeclName()
<< llvm::to_underlying(OldTarget) << OldFD;
Diag(OldFD->getLocation(), diag::note_previous_declaration);
NewFD->setInvalidDecl();
break;
if ((NewTarget == CUDAFunctionTarget::HostDevice &&
!(getLangOpts().OffloadImplicitHostDeviceTemplates &&
isImplicitHostDeviceFunction(NewFD) &&
OldTarget == CUDAFunctionTarget::Device)) ||
(OldTarget == CUDAFunctionTarget::HostDevice &&
!(getLangOpts().OffloadImplicitHostDeviceTemplates &&
isImplicitHostDeviceFunction(OldFD) &&
NewTarget == CUDAFunctionTarget::Device)) ||
(NewTarget == CUDAFunctionTarget::Global) ||
(OldTarget == CUDAFunctionTarget::Global)) {
Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
<< llvm::to_underlying(NewTarget) << NewFD->getDeclName()
<< llvm::to_underlying(OldTarget) << OldFD;
Diag(OldFD->getLocation(), diag::note_previous_declaration);
NewFD->setInvalidDecl();
break;
}
if ((NewTarget == CUDAFunctionTarget::Host &&
OldTarget == CUDAFunctionTarget::Device) ||
(NewTarget == CUDAFunctionTarget::Device &&
OldTarget == CUDAFunctionTarget::Host)) {
Diag(NewFD->getLocation(), diag::warn_offload_incompatible_redeclare)
<< llvm::to_underlying(NewTarget) << llvm::to_underlying(OldTarget);
Diag(OldFD->getLocation(), diag::note_previous_declaration);
}
}
}
}
Expand Down
19 changes: 19 additions & 0 deletions clang/test/SemaCUDA/function-redclare.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 -Wnvcc-compat %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only \
// RUN: -isystem %S/Inputs -fcuda-is-device -Wnvcc-compat -verify=redecl %s

// expected-no-diagnostics
#include "cuda.h"

__device__ void f(); // redecl-note {{previous declaration is here}}

void f() {} // redecl-warning {{target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:new declaration is __host__ function, old declaration is __device__ function}}

void g(); // redecl-note {{previous declaration is here}}

__device__ void g() {} // redecl-warning {{target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:new declaration is __device__ function, old declaration is __host__ function}}
11 changes: 11 additions & 0 deletions llvm/docs/CompileCudaWithLLVM.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -418,6 +418,17 @@ the compiler chooses to inline ``host_only``.
Member functions, including constructors, may be overloaded using H and D
attributes. However, destructors cannot be overloaded.

Clang Warnings for Host and Device Function Declarations
--------------------------------------------------------

Clang can emit warnings when it detects that host (H) and device (D) functions are declared or defined with the same signature. These warnings are not enabled by default.

To enable these warnings, use the following compiler flag:

.. code-block:: console

-Wnvcc-compat

Using a Different Class on Host/Device
--------------------------------------

Expand Down
Loading