[clang] [clang][SYCL] Add sycl_external attribute and restrict emitting device code (PR #140282)

[Tom Honermann via cfe-commits]

================
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
+
+// Semantic tests for sycl_external attribute
+
+[[clang::sycl_external(3)]] // expected-error {{'sycl_external' attribute takes no arguments}}
+void bar() {}
+
+[[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}}
+static void func1() {}
+
+namespace {
+  [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}}
+  void func2() {}
+
+  struct UnnX {};
+}
+
+[[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}}
+  void func4(UnnX) {}
+
+class A {
+  [[clang::sycl_external]]
+  A() {}
+
+  [[clang::sycl_external]] void func3() {}
+};
+
+class B {
+public:
+  [[clang::sycl_external]] virtual void foo() {}
+
+  [[clang::sycl_external]] virtual void bar() = 0;
+};
+
+[[clang::sycl_external]] int *func0() { return nullptr; }
+
+[[clang::sycl_external]] void func2(int *) {}
+
----------------
tahonermann wrote:

The attribute should be usable with (member) function template declarations. Let's exercise those too. See `sycl-kernel-entry-point-attr-appertainment.cpp` for a bunch of other kinds of entities that we could test for rejected appertainment with an appropriate diagnostic.

Other suggested tests:
- `main()` should be rejected.
- Explicitly deleted functions should be rejected (explicitly defaulted functions that are defined as deleted should be accepted).

Finally, we should have a plan to ensure the following diagnostic requirements are satisfied at some point.
- When compiling for a SYCL device target that does not support the generic address space, the function shall not specify a raw pointer or reference type as the return type or as a parameter type.
- The function cannot call `sycl::group::parallel_for_work_item`.
- The function cannot be called from a `parallel_for_work_group` scope.

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