[CUDA][HIP] Fix implicit attribute of builtin

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.

Author: yxsamliu

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

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);
+}