https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/214911
Summary: Currently, if the user redeclares a `__host__ __device__` function it will be assumed to be `__host__` and fail the redeclaration. NVCC accepts this for `__host__ __device__` functions and I believe we should as well. We deliberately ignore the `__device__` only case, this is still an error for us and a warning for NVCC. An example is at https://godbolt.org/z/ahvTcxc6T >From 64307f33061948eeccb5e707c01fa9ea808cea98 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Fri, 7 Aug 2026 20:48:08 -0500 Subject: [PATCH] [CUDA/HIP] Allow __host__ __device__ function redeclarations Summary: Currently, if the user redeclares a `__host__ __device__` function it will be assumed to be `__host__` and fail the redeclaration. NVCC accepts this for `__host__ __device__` functions and I believe we should as well. We deliberately ignore the `__device__` only case, this is still an error for us and a warning for NVCC. An example is at https://godbolt.org/z/ahvTcxc6T --- clang/lib/Sema/SemaDecl.cpp | 12 +++++- clang/lib/Sema/SemaOverload.cpp | 7 ++++ .../CodeGenCUDA/host-device-redeclaration.cu | 41 +++++++++++++++++++ clang/test/SemaCUDA/function-overload.cu | 23 +++++++++++ 4 files changed, 81 insertions(+), 2 deletions(-) create mode 100644 clang/test/CodeGenCUDA/host-device-redeclaration.cu diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index ae8c26808a164..d3988fa1e04fd 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -2963,8 +2963,16 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D, (isa<CUDAHostAttr>(Attr) || isa<CUDADeviceAttr>(Attr) || isa<CUDAGlobalAttr>(Attr))) { // CUDA target attributes are part of function signature for - // overloading purposes and must not be merged. - return false; + // overloading purposes and must not be merged. A redeclaration spelling + // that includes host usage is not a new overload and will be merged. + auto SpelledOnRedecl = [](const auto *A) { + return A && !A->isImplicit() && !A->isInherited(); + }; + if (isa<CUDAGlobalAttr>(Attr) || Attr->isImplicit() || + SpelledOnRedecl(D->getAttr<CUDAHostAttr>()) || + SpelledOnRedecl(D->getAttr<CUDADeviceAttr>())) + return false; + NewAttr = cast<InheritableAttr>(Attr->clone(S.Context)); } else if (const auto *MA = dyn_cast<MinSizeAttr>(Attr)) NewAttr = S.mergeMinSizeAttr(D, *MA); else if (const auto *SNA = dyn_cast<SwiftNameAttr>(Attr)) diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index d66dea0d918fa..8f4e69c7890ea 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -1653,6 +1653,13 @@ static bool IsOverloadOrOverrideImpl(Sema &SemaRef, FunctionDecl *New, !hasExplicitAttr<CUDADeviceAttr>(New)) { return false; } + // Redeclarations inherit behavior for the __host__ __device__ case. + if (OldTarget == CUDAFunctionTarget::HostDevice && + NewTarget == CUDAFunctionTarget::Host && + !hasExplicitAttr<CUDAHostAttr>(New) && + !hasExplicitAttr<CUDADeviceAttr>(New)) { + return false; + } return true; } } diff --git a/clang/test/CodeGenCUDA/host-device-redeclaration.cu b/clang/test/CodeGenCUDA/host-device-redeclaration.cu new file mode 100644 index 0000000000000..db46682ea08bc --- /dev/null +++ b/clang/test/CodeGenCUDA/host-device-redeclaration.cu @@ -0,0 +1,41 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \ +// RUN: | FileCheck --check-prefix=HOST %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \ +// RUN: | FileCheck --check-prefix=DEVICE %s + +#include "Inputs/cuda.h" + +__host__ __device__ int hd(); +int hd() { return 42; } + +struct S { + __host__ __device__ int f(); +}; +int S::f() { return 42; } + +template <typename T> +struct T1 { + __host__ __device__ int f(); +}; +template <typename T> +int T1<T>::f() { return 42; } + +__global__ void kernel(int *out) { + S s; + T1<int> t; + *out = hd() + s.f() + t.f(); +} + +int host_caller() { + S s; + T1<int> t; + return hd() + s.f() + t.f(); +} + +// HOST-DAG: define{{.*}} i32 @_Z2hdv() +// HOST-DAG: define{{.*}} i32 @_ZN1S1fEv( +// HOST-DAG: define{{.*}} i32 @_ZN2T1IiE1fEv( + +// DEVICE-DAG: define{{.*}} i32 @_Z2hdv() +// DEVICE-DAG: define{{.*}} i32 @_ZN1S1fEv( +// DEVICE-DAG: define{{.*}} i32 @_ZN2T1IiE1fEv( diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu index 11f84a912ea7b..0c744bd6d11aa 100644 --- a/clang/test/SemaCUDA/function-overload.cu +++ b/clang/test/SemaCUDA/function-overload.cu @@ -61,6 +61,29 @@ __device__ int dhd() { return 0; } // expected-note {{previous declarat __host__ __device__ int dhd() { return 0; } // expected-error@-1 {{__host__ __device__ function 'dhd' cannot overload __device__ function 'dhd'}} +// A host/device redeclaration with no attribute inherits the initial state. +__host__ __device__ int hd_bare(); +int hd_bare() { return 0; } +__global__ void call_hd_bare() { hd_bare(); } + +struct HDMember { + __host__ __device__ int f(); +}; +int HDMember::f() { return 0; } +__global__ void call_hd_member(HDMember m) { m.f(); } + +template <typename T> +struct HDTemplateMember { + __host__ __device__ int f(); +}; +template <typename T> +int HDTemplateMember<T>::f() { return 0; } +__global__ void call_hd_template_member(HDTemplateMember<int> m) { m.f(); } + +int bare_then_hd(); // expected-note {{previous declaration is here}} +__host__ __device__ int bare_then_hd(); +// expected-error@-1 {{__host__ __device__ function 'bare_then_hd' cannot overload __host__ function 'bare_then_hd'}} + // Same tests for extern "C" functions. extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}} extern "C" int chh() { return 0; } // expected-error {{redefinition of 'chh'}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
