[ROCm OpenCL] device::get_info<device::sub_group_sizes> throws Native API failed

[al42and]

Describe the bug

With ROCm 4.5.2, trying to call device.get_info<sycl::info::device::sub_group_sizes>() on an AMD device throws cl::sycl::runtime_error.

To Reproduce

#include <CL/sycl.hpp>
#include <iostream>

int main() {
  std::vector<sycl::device> devices = sycl::device::get_devices();

  for (const auto &dev : devices) {
    std::cout << "Device " << dev.get_info<sycl::info::device::name>()
              << std::endl;
    const auto sg_sizes = dev.get_info<sycl::info::device::sub_group_sizes>();
  }
}

$ clang++ -g -fsycl sgsize.cpp -o sgsizes && SYCL_DEVICE_FILTER=opencl:gpu ./sgsizes
Device gfx906:sramecc+:xnack-
terminate called after throwing an instance of 'cl::sycl::runtime_error'
  what():  Native API failed. Native API returns: -30 (CL_INVALID_VALUE) -30 (CL_INVALID_VALUE)
Aborted

Works fine with SYCL_DEVICE_FILTER=hip:gpu.

Looking at sycl/plugins/opencl/pi_opencl.cpp, it seems CL_DEVICE_SUB_GROUP_SIZES_INTEL is passed to clGetDeviceInfo, which, being non-Intel, does not know what to do with it.

It is understandable that the OpenCL runtime does not support this extension. But it would be nice if it was handled more gracefully, not with a cryptic Native API failed. Native API returns: -30 (CL_INVALID_VALUE).

Better yet, of course, is to provide the supported sub-group sizes.

Environment (please complete the following information):

• OS: Linux
• Target device and vendor: AMD MI50 (gfx906)
• DPC++ version: 16e1920
• Dependencies version: ROCm 4.5.2

[zjin-lcf]

Is the value of 'sub_group_size' 64 for AMD devices ?

[al42and]

@zjin-lcf AMD RDNA architecture also supports 32-wide execution.

EDIT: But I specifically am more concerned with the function not throwing rather than querying the support for Wave32.

[s-kanaev]

Do I understand correctly, that notion of warps in HIP is the same as in CUDA and it matches the subgroups notion in SYCL?
Then, I reckon this info should use hipDeviceAttributeWarpSize, shouldn't it?

[s-kanaev]

The code is in place for quite a time: https://github.com/intel/llvm/blob/sycl/sycl/plugins/hip/pi_hip.cpp#L1085

@al42and , could you, please, clarify if you encounter exception throw with AMD GPU and HIP backend or OpenCL backend?

@malixian, could you, please, comment here?

[al42and]

Do I understand correctly, that notion of warps in HIP is the same as in CUDA and it matches the subgroups notion in SYCL?

It's called "wavefronts" in AMD lingo, but yes, it's the same thing.

Then, I reckon this info should use hipDeviceAttributeWarpSize, shouldn't it?

Almost. With HIP backend, it indeed should use, and in fact uses, hipDeviceAttributeWarpSize, which returns 64 as expected. But here, the problem is with the OpenCL backend which is happy to use AMD ICD when available:

[opencl:gpu:0] AMD Accelerated Parallel Processing, gfx906:sramecc+:xnack- 2.0 [3361.0 (HSA1.1,LC)]
[opencl:gpu:1] AMD Accelerated Parallel Processing, gfx906:sramecc+:xnack- 2.0 [3361.0 (HSA1.1,LC)]
[ext_oneapi_hip:gpu:0] AMD HIP BACKEND, gfx906:sramecc+:xnack- 0.0 [HIP 40421.43]
[ext_oneapi_hip:gpu:1] AMD HIP BACKEND, gfx906:sramecc+:xnack- 0.0 [HIP 40421.43]

So, the subgroup size detection works fine for ext_oneapi_hip:gpu devices, but crashes for opencl:gpu devices, because AMD's clGetDeviceInfo doesn't support the CL_DEVICE_SUB_GROUP_SIZES_INTEL property.

[al42and]

@al42and , could you, please, clarify if you encounter exception throw with AMD GPU and HIP backend or OpenCL backend?

@s-kanaev, it is with OpenCL backend, as mentioned in the Issue title :)

[s-kanaev]

Sorry, my bad, I overlooked it.

[s-kanaev]

It's not clear how to return sub group size when backend implementation doesn't support it hence, I believe some other exception wording might suite here.

[al42and]

I believe some other exception wording might suite here.

That would already be a big improvement. The current error is generic to the point of uselessness.

But AMD offers CL_DEVICE_WAVEFRONT_WIDTH_AMD property, which should work and would be a better solution.

[rolandschulz]

It's not clear how to return sub group size when backend implementation doesn't support it hence, I believe some other exception wording might suite here.

I don't think that would be OK. The spec doesn't allow for the sub_group_sizes info descriptor to throw.

One possible solution would be to hide any OpenCL devices which don't support the Intel subgroup extension so they can't be queried.