diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9f0b6f5a36389..c272e281c9e8c 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -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< + "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>, 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 80ea43dc5316e..580b9872c6a1d 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -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); + } } } } diff --git a/clang/test/SemaCUDA/function-redclare.cu b/clang/test/SemaCUDA/function-redclare.cu new file mode 100644 index 0000000000000..7cd9bad79ae98 --- /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 -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}} diff --git a/llvm/docs/CompileCudaWithLLVM.rst b/llvm/docs/CompileCudaWithLLVM.rst index 631691ef9b472..0371d7a3bdfcb 100644 --- a/llvm/docs/CompileCudaWithLLVM.rst +++ b/llvm/docs/CompileCudaWithLLVM.rst @@ -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 --------------------------------------