[SYCL] Add property set types and JSON representation

[jzc]

This PR adds the PropertySet type, along with a pair of functions used to serialize and deserialize into a JSON representation. A property set is a key-value map, with values being one of 2 types - uint32 or byte array. A property set registry is a collection of property sets, indexed by a "category" name.

In SYCL offloading, property sets will be used to communicate metadata about device images needed by the SYCL runtime. For example, there is a property set which has a byte array containing the numeric ID, offset, and size of each SYCL2020 spec constant. Another example is a property set describing the optional kernel features used in the module: does it use fp64? fp16? atomic64?

This metadata will be computed by clang-sycl-linker and the JSON representation will be inserted in the string table of each output OffloadBinary. This JSON will be consumed the SYCL offload wrapper and will be lowered to the binary form SYCL runtime expects.

For example, consider this SYCL program that calls a kernel that uses fp64:

#include <sycl/sycl.hpp>

using namespace sycl;
class MyKernel;

int main() {
  queue q;
  auto *p = malloc_shared<double>(1, q);
  *p = .1;
  q.single_task<MyKernel>([=]{ *p *= 2; }).wait();
  std::cout << *p << "\n";
  free(p, q);
}

The device code for this program would have the kernel marked with !sycl_used_aspects:

define spir_kernel void @_ZTS8MyKernel([...]) !sycl_used_aspects !n { [...] }
!n = {i32 6}

clang-sycl-linker would recognize this metadata and then would output the following JSON in the OffloadBinary's key-value map:

{
  "SYCL/device requirements": {
    // aspects contains a list of sycl::aspect values used
    // by the module; in this case just the value 6 encoded 
    // as a 4-byte little-endian integer 
    "aspects": "BjAwMA=="
  }
}

The SYCL offload wrapper would lower those property sets to something like this:

struct _sycl_device_binary_property_set_struct {
  char *CategoryName;
  _sycl_device_binary_property *PropertiesBegin;
  _sycl_device_binary_property *PropertiesEnd;
};

struct _sycl_device_binary_property_struct {
  char *PropertyName;  
  void *ValAddr;     
  uint64_t ValSize; 
};

//  
_sycl_device_binary_property_struct device_requirements[] = {
  /* PropertyName */ "aspects",
  /* ValAddr */ [pointer to the bytes 0x06 0x00 0x00 0x00],
  /* ValSize */ 4,
};

_sycl_device_binary_property_set_struct properties[] = {
  /* CategoryName */ "SYCL/device requirements",
  /* PropertiesBegin */ device_requirements,
  /* PropertiesEnd */ std::end(device_requirments),
}

[maksimsab]

Summary of my opinion.

I think, that the initial design was very poor and it doesn't worth to save it anyhow.

We could come up with the simple design like the following:

namespace sycl {

using PropertyValue = std::variant<uint64_t, std::string>;

using Properties = std::map<std::string, std::map<std::string, PropertyValue>>;

void writePropertiesToJSON(const Properties &P, raw_ostream &O);

Expected<Properties> readPropertiesFromJSON(MemoryBufferRef MBR);

} // namespace sycl

Everything else could be thrown away. And this has the most chances to get upstreamed in my opinion.
Please, share your opinion.

[maksimsab]

Great work in overall!

I have some requests about error messages.

[maksimsab]

Perhaps, we should include missing json header.

[maksimsab]

Maybe, we could use a HEX-encoded string for the case of ByteArray. It would save some space in encoded JSON object because uint64_t is used for each byte and this is a 8x overhead.

[maksimsab]

Any opinion on this one?
This is also possible to reuse base64 encoding that is available in LLVM.

[jzc]

True, I don't mind either way so I switched it to base64

[maksimsab]

Any opinion on this one?
This is also possible to reuse base64 encoding that is available in LLVM.

[maksimsab]

LGTM.

[asudarsa]

Can we not do this?

Suggested change

	void llvm::offloading::writePropertiesToJSON(
	void writePropertiesToJSON(

[asudarsa]

If we can, pls update other places as well...Thanks

[jzc]

It actually seems to be preferred: https://llvm.org/docs/CodingStandards.html#use-namespace-qualifiers-to-implement-previously-declared-functions

[asudarsa]

Please check shared build as well. It is a good idea to do that whenever you are introducing new files (or new headers in files)

Thanks

[asudarsa]

LGTM. Think we can submit this upstream once we have a good PR description with an example and also details on where this will be used
Thanks