Expose buffer page size as property for concurrent kernel buffer access

[illuhad]

This adds the page size property to make the buffer page size user configurable:

namespace sycl::property::buffer {
template<int Dim>
class hipSYCL_page_size
{
public:
  hipSYCL_page_size(const sycl::range<Dim>& page_size);
};
}

In hipSYCL terminology, a page represents a multidimensional chunk of memory (potentially strided) for which the runtime manages data state individually. hipSYCL pages are somewhat inspired by, but unrelated to OS virtual memory pages.
A buffer consists in general of multiple pages. The page size therefore controls the granularity of data management within a buffer.

See here for more information on the hipSYCL memory model:
https://github.com/illuhad/hipSYCL/blob/develop/doc/runtime-spec.md#data-state-tracking-and-pages

In short, this allows the user to depart from the classical SYCL memory model where (apart from two read accesses) no two kernels can access the same buffer concurrently.
In the hipSYCL model, concurrent kernels operating on the same buffer are allowed if the hipSYCL memory pages that are accessed by the two kernels do not intersect.
By default, the page size equals the buffer size to comply with the SYCL spec and forbid concurrent kernels operating on the same buffer. If the user sets a smaller page size using the buffer property, concurrent kernels are possible as long as two kernels don't access the same pages.

This PR therefore allows concurrent kernels operating on the same buffer using an intuitive code pattern like this:

// Construct a buffer consisting of four pages in total
sycl::buffer<int, 2> buff{sycl::range{512, 512},
                          sycl::property::buffer::hipSYCL_page_size<2>{
                              sycl::range{256, 256}}};

// Kernel 1
q.submit([&](sycl::handler &cgh) {
  // accesses page (0, 0) and (0, 1)
  sycl::range range{256, 512};
  sycl::id offset{0, 0};

  sycl::accessor<int, 2> acc{buff, cgh, range, offset};
  cgh.parallel_for(...);
});


// Kernel 2
q.submit([&](sycl::handler &cgh) {
  // accesses page (1, 0) and (1, 1)
  sycl::range range{256, 512};
  sycl::id offset{256, 0};

  sycl::accessor<int, 2> acc{buff, cgh, range, offset};
  cgh.parallel_for(...);
});

In this example, the two kernels have no dependencies on each other in the DAG because they access different pages, so there is no data access conflict.
This can be useful for:

• multiple devices operating on one data buffer
• improved performance on a single device because the runtime is free to execute kernels concurrently

[illuhad]

@hipSYCL-ci-bot test.

[fodinabor]

Looks great :)