Author: Yaxun (Sam) Liu Date: 2025-05-07T22:03:33-04:00 New Revision: c16297cd3f0ed9d036e9cf16fb6885aa3c72d5d3
URL: https://github.com/llvm/llvm-project/commit/c16297cd3f0ed9d036e9cf16fb6885aa3c72d5d3 DIFF: https://github.com/llvm/llvm-project/commit/c16297cd3f0ed9d036e9cf16fb6885aa3c72d5d3.diff LOG: [CUDA][HIP] Fix host/device attribute of builtin (#138162) When a builtin function is passed a pointer with a different address space, clang creates an overloaded builtin function but does not copy the host/device attribute. This causes error when the builtin is called by device functions since CUDA/HIP relies on the host/device attribute to treat a builtin function as callable on both host and device sides. Fixed by copying the host/device attribute of the original builtin function to the created overloaded builtin function. Added: clang/test/SemaCUDA/overloaded-builtin.cu Modified: clang/lib/Sema/SemaExpr.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index c3ef5a70d5f6d..57135adf714ce 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6362,6 +6362,14 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema *Sema, ASTContext &Context, Params.push_back(Parm); } OverloadDecl->setParams(Params); + // We cannot merge host/device attributes of redeclarations. They have to + // be consistent when created. + if (Sema->LangOpts.CUDA) { + if (FDecl->hasAttr<CUDAHostAttr>()) + OverloadDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); + if (FDecl->hasAttr<CUDADeviceAttr>()) + OverloadDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); + } Sema->mergeDeclAttributes(OverloadDecl, FDecl); return OverloadDecl; } diff --git a/clang/test/SemaCUDA/overloaded-builtin.cu b/clang/test/SemaCUDA/overloaded-builtin.cu new file mode 100644 index 0000000000000..c60c27e7f8627 --- /dev/null +++ b/clang/test/SemaCUDA/overloaded-builtin.cu @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -fsyntax-only -verify=host -xhip %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify=dev -xhip %s + +// dev-no-diagnostics + +#include "Inputs/cuda.h" + +__global__ void kernel() { + __attribute__((address_space(0))) void *mem_ptr; + (void)__builtin_amdgcn_is_shared(mem_ptr); +} + +template<typename T> +__global__ void template_kernel(T *p) { + __attribute__((address_space(0))) void *mem_ptr; + (void)__builtin_amdgcn_is_shared(mem_ptr); +} + +void hfun() { + __attribute__((address_space(0))) void *mem_ptr; + (void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}} +} + +template<typename T> +void template_hfun(T *p) { + __attribute__((address_space(0))) void *mem_ptr; + (void)__builtin_amdgcn_is_shared(mem_ptr); // host-error {{reference to __device__ function '__builtin_amdgcn_is_shared' in __host__ function}} +} + + +int main() { + int *p; + kernel<<<1,1>>>(); + template_kernel<<<1,1>>>(p); + template_hfun(p); // host-note {{called by 'main'}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
