llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Yaxun (Sam) Liu (yxsamliu) <details> <summary>Changes</summary> When a builtin function with generic pointer parameter is passed a pointer with address space, clang creates an overloaded builtin function but does not make it implicit. This causes error when the builtin is called by device functions since CUDA/HIP relies on the implicit attribute to treat a builtin function as callable on both host and device sides. Fixed by making the created overloaded builtin functions implicit. --- Full diff: https://github.com/llvm/llvm-project/pull/138162.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaExpr.cpp (+1) - (added) clang/test/SemaCUDA/overloaded-builtin.cu (+23) ``````````diff diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 0cd86dc54b0ab..d9eccb31e6d1e 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6358,6 +6358,7 @@ static FunctionDecl *rewriteBuiltinFunctionDecl(Sema *Sema, ASTContext &Context, } OverloadDecl->setParams(Params); Sema->mergeDeclAttributes(OverloadDecl, FDecl); + OverloadDecl->setImplicit(true); return OverloadDecl; } diff --git a/clang/test/SemaCUDA/overloaded-builtin.cu b/clang/test/SemaCUDA/overloaded-builtin.cu new file mode 100644 index 0000000000000..719bfea4aef2f --- /dev/null +++ b/clang/test/SemaCUDA/overloaded-builtin.cu @@ -0,0 +1,23 @@ +// expected-no-diagnostics + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -fsyntax-only -verify -xhip %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsyntax-only -fcuda-is-device -verify -xhip %s + +#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); +} + +int main() { + int *p; + kernel<<<1,1>>>(); + template_kernel<<<1,1>>>(p); +} `````````` </details> https://github.com/llvm/llvm-project/pull/138162 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
