|
| 1 | +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -std=c++17 -sycl-std=2020 -verify -fsyntax-only %s |
| 2 | + |
| 3 | +#include "sycl.hpp" |
| 4 | + |
| 5 | +// This test shows that global_variable_allowed attribute allows |
| 6 | +// global variables of type decorated with it to be referenced in device code. |
| 7 | + |
| 8 | +template <typename T> |
| 9 | +struct [[__sycl_detail__::global_variable_allowed]] global_variable_allowed { |
| 10 | +public: |
| 11 | + const T &get() const noexcept { return *Data; } |
| 12 | + global_variable_allowed() {} |
| 13 | + operator T &() noexcept { return *Data; } |
| 14 | + |
| 15 | +private: |
| 16 | + T *Data; |
| 17 | +}; |
| 18 | + |
| 19 | +global_variable_allowed<int> glob; |
| 20 | +static global_variable_allowed<float> static_glob; |
| 21 | +inline global_variable_allowed<double> inline_glob; |
| 22 | +static const global_variable_allowed<int> static_const_glob; |
| 23 | + |
| 24 | +struct Foo { |
| 25 | + static global_variable_allowed<char> d; |
| 26 | +}; |
| 27 | +global_variable_allowed<char> Foo::d; |
| 28 | +global_variable_allowed<Foo> foo_instance; |
| 29 | +static global_variable_allowed<Foo> foo_static_instance; |
| 30 | + |
| 31 | +extern global_variable_allowed<int> ext; |
| 32 | + |
| 33 | +template <typename T> |
| 34 | +struct Bar { |
| 35 | + static global_variable_allowed<T> e; |
| 36 | +}; |
| 37 | +static const Bar<int> f; |
| 38 | + |
| 39 | +namespace baz { |
| 40 | + global_variable_allowed<int> baz_glob; |
| 41 | +} |
| 42 | + |
| 43 | +// expected-error@+1{{'global_variable_allowed' attribute only applies to classes}} |
| 44 | +[[__sycl_detail__::global_variable_allowed]] int integer; |
| 45 | + |
| 46 | +// expected-error@+1{{'global_variable_allowed' attribute only applies to classes}} |
| 47 | +[[__sycl_detail__::global_variable_allowed]] int *pointer; |
| 48 | + |
| 49 | +// expected-error@+1{{'global_variable_allowed' attribute takes no arguments}} |
| 50 | +struct [[__sycl_detail__::global_variable_allowed(72)]] attribute_argument; |
| 51 | + |
| 52 | + |
| 53 | +int main() { |
| 54 | + sycl::queue q; |
| 55 | + q.submit([&](sycl::handler &h) { |
| 56 | + h.single_task<class KernelName1>([=]() { |
| 57 | + (void)glob; |
| 58 | + (void)static_glob; |
| 59 | + (void)inline_glob; |
| 60 | + (void)static_const_glob; |
| 61 | + (void)Foo::d; |
| 62 | + (void)foo_instance; |
| 63 | + (void)foo_static_instance; |
| 64 | + (void)ext; |
| 65 | + (void)f; |
| 66 | + }); |
| 67 | + }); |
| 68 | +} |
0 commit comments