================
@@ -0,0 +1,109 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang -xhip --offload-arch=gfx1030 --offload-host-only -pedantic
-nogpuinc -nogpulib -nobuiltininc -fsyntax-only -Xclang -verify=onhost %s
+// RUN: %clang -xhip --offload-arch=gfx1030 --offload-device-only -pedantic
-nogpuinc -nogpulib -nobuiltininc -fsyntax-only -Xclang -verify=ondevice %s
+
+// ondevice-no-diagnostics
+
+#include <type_traits>
+
+#define WRAPPED __AMDGCN_WAVEFRONT_SIZE__
+
+#define DOUBLE_WRAPPED (WRAPPED)
+
+__attribute__((host, device)) void use(int, const char*);
+
+template<int N> __attribute__((host, device)) int templatify(int x) {
+ return x + N;
+}
+
+// no warning expected
+#if defined(__HIP_DEVICE_COMPILE__) && (__AMDGCN_WAVEFRONT_SIZE__ == 64) &&
(__AMDGCN_WAVEFRONT_SIZE == 64)
+int foo(void);
+#endif
+
+// no warning expected
+__attribute__((device)) int device_var = __AMDGCN_WAVEFRONT_SIZE__;
+
+__attribute__((device))
+void device_fun() {
+ // no warnings expected
+ use(__AMDGCN_WAVEFRONT_SIZE, "device function");
+ use(__AMDGCN_WAVEFRONT_SIZE__, "device function");
+ use(WRAPPED, "device function");
+ use(DOUBLE_WRAPPED, "device function");
+ use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "device function");
+}
+
+__attribute__((global))
+void global_fun() {
+ // no warnings expected
+ use(__AMDGCN_WAVEFRONT_SIZE, "global function");
+ use(__AMDGCN_WAVEFRONT_SIZE__, "global function");
+ use(WRAPPED, "global function");
+ use(DOUBLE_WRAPPED, "global function");
+ use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "global function");
+}
+
+// warning expected
+int host_var = __AMDGCN_WAVEFRONT_SIZE__; // onhost-warning {{device-specific
macro __AMDGCN_WAVEFRONT_SIZE__ is not available in a __host__ context}}
----------------
Artem-B wrote:
Will `const` or `constexpr` host variables dependent on the macros also produce
warnings?
E.g. something like this https://godbolt.org/z/1bxnrxrnn may be OK:
```
const int z = __AMDGCN_WAVEFRONT_SIZE;
__global__ void kernel(int* array, int n) {
do_something_with(z);
}
```
The use of `z` on the host side would still be wrong, though.
https://github.com/llvm/llvm-project/pull/109663
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
