================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
