[SYCL] Add basic support for the generic_space address space

[ghost]

The basic support is added to let the sycl::atomic_ref class template
be instantiated with the address_space::generic_space address space.

[ghost]

I'm working on a patch for the llvm-test-suite project to test sycl::atomic_refs with generic_space as well.

[vladimirlaz]

/summary:run

[zjin-lcf]

The basic support is added to let the sycl::atomic_ref class template be instantiated with the address_space::generic_space address space.

I have a question. When should generic_space be used in an example/kernel ? Thanks

[ghost]

@zjin-lcf I've found only the sentences about why generic address space can be not the best choice from the performance and linking point of view in the beginning of Section 4.7.7 Address space classes in the SYCL 2020 specification. I hope @bader or @gmlueck can answer more detailed.

[Naghasan]

When should generic_space be used in an example/kernel ?

it was added for completeness rather than actual usefulness. Without you would force users to specialize for raw pointers.

that accounts for that kind of use case:

template<access::address_space Space>
void bar(multi_ptr<int, Space>);

void foo(int* p) {
  // without access::address_space::generic_space, you need to specialize bar or
  // try to find out to which address space p is pointing to.
  bar<access::address_space::generic_space>(p); 
}

[zjin-lcf]

Thank you for the example and comments.

[vladimirlaz]

/summary:run

[ghost]

/summary:run

[ghost]

Please have a look at PR intel/llvm-test-suite#619, the tests against this PR failed but not in the AtomicRef directory.