[SYCL][CUDA] Allow varying program metadata in CUDA backend

[steffenlarsen]

Recently the program metadata produced for reqd_work_group_size changed to allow between 1 and 3 values instead of guaranteing 3 values. This commit makes the CUDA backend correctly parse reqd_work_group_size metadata with variable number of elements.

[steffenlarsen]

Currently in draft as sycl::reqd_work_group_size do not currently generate the metadata needed for sycl-post-link to produce program metadata with < 3 values. See #7450.
Additionally, there is a known bug with the CUDA backend not correctly handling the metadata produced for the work_group_size kernel property, so these cannot currently be used to test this change.

[zjin-lcf]

Can these changes be applied to the HIP backend ?

[steffenlarsen]

Can these changes be applied to the HIP backend ?

To my knowledge the HIP backend doesn't currently handle this metadata. In the future I would like to move this logic out to the runtime level, but we are not quite ready for it yet. However, as the HIP backend doesn't support it currently I figure unless we have a good reason to make the throw-away work it will be better to wait.

[zjin-lcf]

Since the metadata (i.e. reqd_work_group_size) is available in a device binary, could you please explain why the metadata is not available in the HIP backend ? Thanks.

TEST_F(CudaKernelsTest, PICreateProgramAndKernelWithMetadata) {

  std::vector<uint32_t> reqdWorkGroupSizeMD;
  reqdWorkGroupSizeMD.reserve(5);
  // 64-bit representing bit size
  reqdWorkGroupSizeMD.push_back(96);
  reqdWorkGroupSizeMD.push_back(0);
  // reqd_work_group_size x
  reqdWorkGroupSizeMD.push_back(8);
  // reqd_work_group_size y
  reqdWorkGroupSizeMD.push_back(16);
  // reqd_work_group_size z
  reqdWorkGroupSizeMD.push_back(32);

  const char *reqdWorkGroupSizeMDConstName =
      "_Z8myKernelPi@reqd_work_group_size";
...

[steffenlarsen]

Since the metadata (i.e. reqd_work_group_size) is available in a device binary, could you please explain why the metadata is not available in the HIP backend ? Thanks.

TEST_F(CudaKernelsTest, PICreateProgramAndKernelWithMetadata) {

  std::vector<uint32_t> reqdWorkGroupSizeMD;
  reqdWorkGroupSizeMD.reserve(5);
  // 64-bit representing bit size
  reqdWorkGroupSizeMD.push_back(96);
  reqdWorkGroupSizeMD.push_back(0);
  // reqd_work_group_size x
  reqdWorkGroupSizeMD.push_back(8);
  // reqd_work_group_size y
  reqdWorkGroupSizeMD.push_back(16);
  // reqd_work_group_size z
  reqdWorkGroupSizeMD.push_back(32);

  const char *reqdWorkGroupSizeMDConstName =
      "_Z8myKernelPi@reqd_work_group_size";
...

Hi @zjin-lcf ! Currently the program metadata is only generated for the PTX target as it was the only target needing it back when the metadata was added. It should be possible to enable it for the HIP backend too, but we are planning on removing the need for the reqd_work_group_size program metadata since the information is made available through kernel information the runtime now receives.