[clang] [CUDA][HIP] Fix implicit attribute of builtin (PR #138162)

[Artem Belevich via cfe-commits]

================
@@ -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);
----------------
Artem-B wrote:

Are these builtins GPU-only?

Perhaps the root cause here is that the builtins do not have the `__device__` attribute, and making them implicit just happens to work around that. If the builtin is not intended to be used in the host code, `__device__` may be what we want to be set on all builtins like that. Not sure if we currently have a way to express that in the builtin definitions in clang.

https://github.com/llvm/llvm-project/pull/138162