[SYCL][HIP] Fix hip prefetch/mem_advise (noop) for rocm < 5.0 (#10958)

The main fix here is to enable prefetch functionality when the HIP
backend is built with rocm versions < 5.
The change to the prefix/mem_advise UR hip adapter APIs fixes four
e2e-tests that fail test-e2e only for HIP 4.x versions (These failures
don't come up in the CI since it tests using rocm 5.x).

The change to Tracing/image_printers.cpp is effectively re-XFailing this
test for rocm 4 only. rocm 4 doesn't support `hipCreateSurfaceObject`
which is called by this test. I'm not sure this legacy image
functionality (Which as I understand it will be replaced by bindless
images) is really working for rocm 5 in a meaningful way, since most
other legacy image e2e tests are XFAIL for hip. But this test can still
be useful for us for ROCM 5 compatibility testing. The XFAIL was
recently removed here
745febe
which led to us finding the rocm 4 fail. I guess that the CI at some
point switched from testing rocm 4 to testing rocm 5, which meant this
test stopped failing in CI.

---------

Signed-off-by: Jack Kirk <jack.kirk@codeplay.com>

Author: JackAKirk

sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp

@@ -1315,7 +1315,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(
     ur_queue_handle_t hQueue, const void *pMem, size_t size,
     ur_usm_migration_flags_t flags, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
-#if HIP_VERSION_MAJOR >= 5
   void *HIPDevicePtr = const_cast<void *>(pMem);
   ur_device_handle_t Device = hQueue->getContext()->getDevice();
 
@@ -1342,12 +1341,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(
     return UR_RESULT_ERROR_ADAPTER_SPECIFIC;
   }
 
+  // HIP_POINTER_ATTRIBUTE_RANGE_SIZE is not an attribute in ROCM < 5,
+  // so we can't perform this check for such cases.
+#if HIP_VERSION_MAJOR >= 5
   unsigned int PointerRangeSize = 0;
   UR_CHECK_ERROR(hipPointerGetAttribute(&PointerRangeSize,
                                         HIP_POINTER_ATTRIBUTE_RANGE_SIZE,
                                         (hipDeviceptr_t)HIPDevicePtr));
   UR_ASSERT(size <= PointerRangeSize, UR_RESULT_ERROR_INVALID_SIZE);
-
+#endif
   // flags is currently unused so fail if set
   if (flags != 0)
     return UR_RESULT_ERROR_INVALID_VALUE;
@@ -1376,28 +1378,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(
   }
 
   return Result;
-#else
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
-#endif
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL
 urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size,
                    ur_usm_advice_flags_t, ur_event_handle_t *phEvent) {
-#if HIP_VERSION_MAJOR >= 5
   void *HIPDevicePtr = const_cast<void *>(pMem);
+// HIP_POINTER_ATTRIBUTE_RANGE_SIZE is not an attribute in ROCM < 5,
+// so we can't perform this check for such cases.
+#if HIP_VERSION_MAJOR >= 5
   unsigned int PointerRangeSize = 0;
   UR_CHECK_ERROR(hipPointerGetAttribute(&PointerRangeSize,
                                         HIP_POINTER_ATTRIBUTE_RANGE_SIZE,
                                         (hipDeviceptr_t)HIPDevicePtr));
   UR_ASSERT(size <= PointerRangeSize, UR_RESULT_ERROR_INVALID_SIZE);
-
+#endif
   // TODO implement a mapping to hipMemAdvise once the expected behaviour
   // of urEnqueueUSMAdvise is detailed in the USM extension
   return urEnqueueEventsWait(hQueue, 0, nullptr, phEvent);
-#else
-  return UR_RESULT_ERROR_UNSUPPORTED_FEATURE;
-#endif
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill2D(

sycl/test-e2e/Tracing/image_printers.cpp

@@ -29,11 +29,15 @@ int main() {
   {
     sycl::image<2> Img(ImgHostData.data(), ChanOrder, ChanType, ImgSize);
     queue Q;
+
+// legacy Images uses an API that is not supported in hip 4.x
+#if HIP_VERSION_MAJOR >= 5
     Q.submit([&](sycl::handler &CGH) {
       auto ImgAcc = Img.get_access<sycl::float4, SYCLWrite>(CGH);
 
       CGH.single_task<class EmptyTask>([=]() {});
     });
+#endif
   }
   return 0;
 }