[clang] [Clang][AMDGPU] Add a new builtin type for buffer rsrc (PR #94830)

[Yaxun Liu via cfe-commits]

================
@@ -0,0 +1,6 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -aux-triple amdgcn-amd-amdhsa %s -fsyntax-only -verify
+
+#define __device__ __attribute__((device))
+
+__device__ __amdgcn_buffer_rsrc_t test_buffer_rsrc_t_device() {} // expected-warning {{non-void function does not return a value}}
+__amdgcn_buffer_rsrc_t test_buffer_rsrc_t_host() {} // expected-error {{'__amdgcn_buffer_rsrc_t' can only be used in device-side function}}
----------------
yxsamliu wrote:

As discussed in https://github.com/llvm/llvm-project/pull/69366, I think the trend is to make HIP more like C++ where every function is both device and host function, and de-emphasize handling based on host/device attributes. Ideally, we can imagine we are compiling a HIP program for a processor that has the capability of both the host CPU and the device GPU, so that we can ignore host/device difference during semantic checking, and we defer the diagnosing to codegen or linker.

The reason is that C++ is not designed with host/device in mind and the current parser/sema does not consider host/device attributes in many cases, especially about templates. Adding more host/device based sema seems to make things more complicated and not to help making generic C++ code (e.g. the standard C++ library) work for both host/device. Another reason not to emphasize the host/device difference is that difference in device/host AST risks violation of ODR and causes issues difficult to diagnose.

In a word, I would not recommend restricting a type to device only. 

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