Add device info query to report support for native asserts.

[aarongreig]

This allows cuda and hip to stop reporting the opencl extension string, see issue #1374

[aarongreig]

LLVM PR intel/llvm#15929

[JackAKirk]

HIP supports native asserts of devices that are most relevant for the users of dpc++ AFAIK: the assert tests definitely pass on CDNA2 cards, and AMD devs implied they should be supported on all supported devices. There was a failure on a RDNA2 card but that has little relevance to dpc++ users.

[aarongreig]

they still have UNSUPPORTED: hip in the e2e test along with a tracker issue, I could change this to a note about it not necessarily being 100% supported and link the issue, or I'm happy to just delete the comment if you don't think it's useful to track that here

[JackAKirk]

they still have UNSUPPORTED: hip in the e2e test along with a tracker issue, I could change this to a note about it not necessarily being 100% supported and link the issue, or I'm happy to just delete the comment if you don't think it's useful to track that here

Sure, whatever you think is best. The only issue with the current comment is that it is all out of date/ factually incorrect.

As far as hip tests marked unsupported in intel/llvm, for a lot of cases this is because the intel/llvm CI does not use an officially supported rocm card. This case is a bit different because the failure still does occur on some officially supported cards in addition to the unsupported ones but I think your wish to add a comment to state something along the lines of "this is mostly supported" is fine.

It is best to try not to take too much account of such failures, unfortunately we have to rely on other testing.

[aarongreig]

I decided to omit the comment, looking at intel/llvm#7634 it does seem like the remaining fail there is a weird issue with how the hardware is set up, and based on Georgi's comment it sounds like confidence is pretty high that this should be upgraded to an unqualified "supported"

[GeorgeWeb]

@aarongreig According to my latest testing a CDNA2 AMD GPU had reliable native asserts support with latest versions of ROCm. However, I am not sure how well things are with consumer grade ones, but I suspect all supported cards should in theory have support for the feature these days.

Just a question/thought on how this ties to DPC++, but not a requirement:
I am wondering whether we need an intel/llvm PR for this too to update the device config file here https://github.com/intel/llvm/blob/d13019638e4ed3078a52cf06a91080875d2e0d0e/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td#L293 wrt native_assert aspect support for HIP. This Device Configuration File is used for SYCL/DPC++ all_devices_have<sycl::aspect> and any_device_has<sycl::aspect> compile-time aspect traits implementation (see https://intel.github.io/llvm/design/DeviceAspectTraitDesign.html). Maybe we can selectively add AspectExt_oneapi_native_assert to all of the Officially supported: HIP-AMD targets regardless of whether they are CDNA arches or whether we tested them all (because it's impossible for us to fully test that on all gpus so we ought to rely on ROCm actually supporting their supported hardware). Until now I've only added it to gfx90a because that's the only one I knew of passing the test reliably 🤷‍♂️.

[JackAKirk]

LGTM for cuda/hip

[aarongreig]

@aarongreig According to my latest testing a CDNA2 AMD GPU had reliable native asserts support with latest versions of ROCm. However, I am not sure how well things are with consumer grade ones, but I suspect all supported cards should in theory have support for the feature these days.

Just a question/thought for this ties to DPC++, but not a requirement: I am wondering whether we need an intel/llvm PR for this too to update the device config file here https://github.com/intel/llvm/blob/d13019638e4ed3078a52cf06a91080875d2e0d0e/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td#L293 wrt native_assert aspect support for HIP. This Device Configuration File is used for SYCL/DPC++ all_devices_have<sycl::aspect> and any_device_has<sycl::aspect> compile-time aspect traits implementation (see https://intel.github.io/llvm/design/DeviceAspectTraitDesign.html). Maybe we can selectively add AspectExt_oneapi_native_assert to all of the Officially supported: HIP-AMD targets regardless of whether they are CDNA arches or whether we tested them all (because it's impossible for us to fully test that on all gpus so we ought to rely on ROCm actually supporting their supported hardware). Until now I've only added it to gfx90a because that's the only one I knew of passing the test reliably 🤷‍♂️.

That probably does need to be updated, although I don't really want to include that in intel/llvm#15929 since it's supposed to just be updating how the internal reporting mechanism works

[GeorgeWeb]

That probably does need to be updated, although I don't really want to include that in intel/llvm#15929 since it's supposed to just be updating how the internal reporting mechanism works

Yeah, I agree. Guess the comment was more of a note to refer to in a future follow-up to that specific part of intel/llvm.

Regardless, all cuda/hip related changes look good! Cheers.

[nrspruit]

LGTM for L0

[aarongreig]

ping @oneapi-src/unified-runtime-native-cpu-write @oneapi-src/unified-runtime-opencl-write