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