Scalar overloads of copy functions

[mkinsner]

@keryell noted in a review that it would be useful to have scalar overloads of the queue and handler copy functions. This likely makes sense and applies to SYCL broadly, so tracking here.

[gmlueck]

Capturing some information here in case the link to the review noted above goes away. The original comment was about this API, which is a SYCL extension in DPC++:

class queue {
public:
  // Copy to device_global
  template <typename T, typename PropertyListT>
  event copy(const std::remove_all_extents_t<T> *src,
    device_global<T, PropertyListT>& dest,
    size_t count = sizeof(T) / sizeof(std::remove_all_extents_t<T>),
    size_t startIndex = 0);

A typical use would be like:

sycl::ext::oneapi::device_global<int> dev_var;

void func(sycl::queue q) {
  int val = 42;
  q.copy(&val, dev_var).wait();  // Write value 42 to "dev_var"

  q.submit([&](sycl::handler &cgh) {
    cgh.single_task([=] {
      int i = dev_var;   // Kernel can now read 42 from "dev_var"
    });
  });
}

And @keryell's comment was:

Perhaps we could have an overload in SYCL for scalars too so it is possible to avoid the &?

Where he was referring to the &val argument to q.copy() in the example use case.

We opened this issue in the Khronos database because we think a similar argument exists for the queue::copy() functions in the core SYCL spec. Should there be shortcut functions for those if the type T is a scalar?

[gmlueck]

I'm not sure if this is what @keryell had in mind, but we could consider making two changes for these shortcut functions:

• Avoid passing a pointer to a host scalar value.
• Make the API synchronous.

This might simplify code for a common case where the host reads or writes a scalar value from a USM pointer. The idea would also extend to the device_global extension above. These shortcut functions would probably read more naturally if they were named read() and write(). For example:

sycl::ext::oneapi::device_global<int> dev_var;

void func(sycl::queue q) {
  int *pval = malloc_device<int>(1, q);  // Allocate an "int" in USM device memory

  q.write(42, pval);    // Write value 42 to the USM device buffer
  q.write(43, dev_var); // Write value 43 to "dev_var"

  q.submit([&](sycl::handler &cgh) {
    cgh.single_task([=] {
      int tmp = dev_var;
      dev_var = *pval;
      *pval = tmp;
    });
  }).wait();

  int i = q.read(pval);    // Read value 43 from USM device buffer
  int j = q.read(dev_var); // Read value 42 from "dev_var"
}

[keryell]

Thank you folks for creating issues for me. :-)
Yes, that looks quite cleaner!

[fraggamuffin]

Not everyone needs to implement device global extension
its a synchronous function with submit or handler so need a separate api, may be free function?