From aa4b7527222522168147c367d614ffbe589fce0d Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Thu, 20 Jul 2023 10:23:17 +0100 Subject: [PATCH 01/17] [SYCL][HIP] Initial HIP mem_advise implementation. --- sycl/plugins/unified_runtime/pi2ur.hpp | 12 ++ .../ur/adapters/hip/enqueue.cpp | 141 ++++++++++++++++-- sycl/test-e2e/USM/memadvise_cuda.cpp | 3 +- 3 files changed, 144 insertions(+), 12 deletions(-) diff --git a/sycl/plugins/unified_runtime/pi2ur.hpp b/sycl/plugins/unified_runtime/pi2ur.hpp index 83f8525054523..4cb0fadddf783 100644 --- a/sycl/plugins/unified_runtime/pi2ur.hpp +++ b/sycl/plugins/unified_runtime/pi2ur.hpp @@ -3228,6 +3228,18 @@ inline pi_result piextUSMEnqueueMemAdvise(pi_queue Queue, const void *Ptr, if (Advice & PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION) { UrAdvice |= UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION; } + if (Advice & PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY) { + UrAdvice |= UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE; + } + if (Advice & PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY) { + UrAdvice |= UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE; + } + if (Advice & PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST) { + UrAdvice |= UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_HOST; + } + if (Advice & PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST) { + UrAdvice |= UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_HOST; + } if (Advice & PI_MEM_ADVICE_RESET) { UrAdvice |= UR_USM_ADVICE_FLAG_DEFAULT; } diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp index 1b0b2acc2a3f8..1e1dc4a7a7c51 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp @@ -82,6 +82,61 @@ void simpleGuessLocalWorkSize(size_t *ThreadsPerBlock, --ThreadsPerBlock[0]; } } + +ur_result_t setHipMemAdvise(const void *DevPtr, size_t Size, + ur_usm_advice_flags_t URAdviceFlags, + hipDevice_t Device) { + std::unordered_map + URToHIPMemAdviseDeviceFlagsMap = { + {UR_USM_ADVICE_FLAG_SET_READ_MOSTLY, hipMemAdviseSetReadMostly}, + {UR_USM_ADVICE_FLAG_CLEAR_READ_MOSTLY, hipMemAdviseUnsetReadMostly}, + {UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION, + hipMemAdviseSetPreferredLocation}, + {UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION, + hipMemAdviseUnsetPreferredLocation}, + {UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE, + hipMemAdviseSetAccessedBy}, + {UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE, + hipMemAdviseUnsetAccessedBy}, + }; + for (auto &FlagPair : URToHIPMemAdviseDeviceFlagsMap) { + if (URAdviceFlags & FlagPair.first) { + UR_CHECK_ERROR(hipMemAdvise(DevPtr, Size, FlagPair.second, Device)); + } + } + + static std::unordered_map + URToHIPMemAdviseHostFlagsMap = { + {UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION_HOST, + hipMemAdviseSetPreferredLocation}, + {UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST, + hipMemAdviseUnsetPreferredLocation}, + {UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_HOST, hipMemAdviseSetAccessedBy}, + {UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_HOST, + hipMemAdviseUnsetAccessedBy}, + }; + + for (auto &FlagPair : URToHIPMemAdviseHostFlagsMap) { + if (URAdviceFlags & FlagPair.first) { + UR_CHECK_ERROR( + hipMemAdvise(DevPtr, Size, FlagPair.second, hipCpuDeviceId)); + } + } + + static constexpr std::array UnmappedMemAdviceFlags = + {UR_USM_ADVICE_FLAG_SET_NON_ATOMIC_MOSTLY, + UR_USM_ADVICE_FLAG_CLEAR_NON_ATOMIC_MOSTLY, + UR_USM_ADVICE_FLAG_BIAS_CACHED, UR_USM_ADVICE_FLAG_BIAS_UNCACHED}; + + for (auto &UnmappedFlag : UnmappedMemAdviceFlags) { + if (URAdviceFlags & UnmappedFlag) { + return UR_RESULT_ERROR_INVALID_ENUMERATION; + } + } + + return UR_RESULT_SUCCESS; +} + } // namespace UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferWrite( @@ -1328,23 +1383,87 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch( #endif } +/// USM: memadvise API to govern behavior of automatic migration mechanisms 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 + ur_usm_advice_flags_t advice, ur_event_handle_t *phEvent) { + UR_ASSERT(pMem && size > 0, UR_RESULT_ERROR_INVALID_VALUE); void *HIPDevicePtr = const_cast(pMem); + + // Passing MEM_ADVISE_SET/MEM_ADVISE_CLEAR_PREFERRED_LOCATION and + // to hipMemAdvise on a GPU device requires the GPU device to report a + // non-zero value for hipDeviceAttributeConcurrentManagedAccess. Therfore, + // ignore memory advise if concurrent managed memory access is not available. + if ((advice & UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION) || + (advice & UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION) || + (advice & UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE) || + (advice & UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE) || + (advice & UR_USM_ADVICE_FLAG_DEFAULT)) { + ur_device_handle_t Device = hQueue->getContext()->getDevice(); + if (!getAttribute(Device, hipDeviceAttributeConcurrentManagedAccess)) { + setErrorMessage("mem_advise ignored as device does not support " + "concurrent managed access", + UR_RESULT_SUCCESS); + return UR_RESULT_ERROR_ADAPTER_SPECIFIC; + } + + // TODO: If pMem points to valid system-allocated pageable memory, we should + // check that the device also has the hipDeviceAttributePageableMemoryAccess + // property. + } + unsigned int PointerRangeSize = 0; - UR_CHECK_ERROR(hipPointerGetAttribute(&PointerRangeSize, - HIP_POINTER_ATTRIBUTE_RANGE_SIZE, - (hipDeviceptr_t)HIPDevicePtr)); + UR_CHECK_ERROR(hipPointerGetAttribute( + &PointerRangeSize, HIP_POINTER_ATTRIBUTE_RANGE_SIZE, + static_cast(HIPDevicePtr))); UR_ASSERT(size <= PointerRangeSize, UR_RESULT_ERROR_INVALID_SIZE); - // 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_result_t Result = UR_RESULT_SUCCESS; + std::unique_ptr EventPtr{nullptr}; + + try { + ScopedContext Active(hQueue->getDevice()); + + if (phEvent) { + EventPtr = + std::unique_ptr(ur_event_handle_t_::makeNative( + UR_COMMAND_USM_ADVISE, hQueue, hQueue->getNextTransferStream())); + EventPtr->start(); + } + + if (advice & UR_USM_ADVICE_FLAG_DEFAULT) { + UR_CHECK_ERROR(hipMemAdvise(pMem, size, hipMemAdviseUnsetReadMostly, + hQueue->getContext()->getDevice()->get())); + UR_CHECK_ERROR(hipMemAdvise(pMem, size, + hipMemAdviseUnsetPreferredLocation, + hQueue->getContext()->getDevice()->get())); + UR_CHECK_ERROR(hipMemAdvise(pMem, size, hipMemAdviseUnsetAccessedBy, + hQueue->getContext()->getDevice()->get())); + } else { + Result = setHipMemAdvise(HIPDevicePtr, size, advice, + hQueue->getContext()->getDevice()->get()); + // UR_RESULT_ERROR_INVALID_ENUMERATION is returned when using a valid but + // currently unmapped advice arguments as not supported by this platform. + // Therefore, warn the user instead of throwing and aborting the runtime. + if (Result == UR_RESULT_ERROR_INVALID_ENUMERATION) { + setErrorMessage("mem_advise is ignored as the advice argument is not " + " supported by this device.", + Result); + return UR_RESULT_ERROR_ADAPTER_SPECIFIC; + } + } + + if (phEvent) { + Result = EventPtr->record(); + *phEvent = EventPtr.release(); + } + } catch (ur_result_t err) { + Result = err; + } catch (...) { + Result = UR_RESULT_ERROR_UNKNOWN; + } + + return Result; } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill2D( diff --git a/sycl/test-e2e/USM/memadvise_cuda.cpp b/sycl/test-e2e/USM/memadvise_cuda.cpp index cacd5be1ae7bf..629cc258b3c22 100755 --- a/sycl/test-e2e/USM/memadvise_cuda.cpp +++ b/sycl/test-e2e/USM/memadvise_cuda.cpp @@ -1,5 +1,5 @@ // RUN: %{build} -o %t1.out -// REQUIRES: cuda +// REQUIRES: cuda || hip_amd // RUN: %{run} %t1.out //==---------------- memadvise_cuda.cpp ------------------------------------==// @@ -32,6 +32,7 @@ int main() { return -1; } + // NOTE: PI_MEM_ADVICE_CUDA_* advice values are mapped to the HIP backend too. std::vector valid_advices{ PI_MEM_ADVICE_CUDA_SET_READ_MOSTLY, PI_MEM_ADVICE_CUDA_UNSET_READ_MOSTLY, From c4869846a4f9d065c51d74b663b8fa3e67115cd5 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Mon, 7 Aug 2023 10:47:13 +0100 Subject: [PATCH 02/17] Add back the HIP version guard for hipMemAdvise. --- sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp index 1e1dc4a7a7c51..cff80f6f5de0c 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp @@ -1387,6 +1387,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch( UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, ur_usm_advice_flags_t advice, ur_event_handle_t *phEvent) { +#if HIP_VERSION_MAJOR >= 5 UR_ASSERT(pMem && size > 0, UR_RESULT_ERROR_INVALID_VALUE); void *HIPDevicePtr = const_cast(pMem); @@ -1464,6 +1465,9 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, } return Result; +#else + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif } UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill2D( From f714d74ecc6396622ad48a0025fd2a6b737202bd Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 8 Aug 2023 15:29:16 +0100 Subject: [PATCH 03/17] Address review simplifying impl and add comments. --- .../ur/adapters/hip/enqueue.cpp | 114 +++++++++--------- 1 file changed, 58 insertions(+), 56 deletions(-) diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp index cff80f6f5de0c..00c8e63294ed6 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp @@ -86,52 +86,53 @@ void simpleGuessLocalWorkSize(size_t *ThreadsPerBlock, ur_result_t setHipMemAdvise(const void *DevPtr, size_t Size, ur_usm_advice_flags_t URAdviceFlags, hipDevice_t Device) { - std::unordered_map - URToHIPMemAdviseDeviceFlagsMap = { - {UR_USM_ADVICE_FLAG_SET_READ_MOSTLY, hipMemAdviseSetReadMostly}, - {UR_USM_ADVICE_FLAG_CLEAR_READ_MOSTLY, hipMemAdviseUnsetReadMostly}, - {UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION, - hipMemAdviseSetPreferredLocation}, - {UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION, - hipMemAdviseUnsetPreferredLocation}, - {UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE, - hipMemAdviseSetAccessedBy}, - {UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE, - hipMemAdviseUnsetAccessedBy}, + using ur_to_hip_advice_t = std::pair; + + static constexpr std::array + URToHIPMemAdviseDeviceFlags{ + std::make_pair(UR_USM_ADVICE_FLAG_SET_READ_MOSTLY, + hipMemAdviseSetReadMostly), + std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_READ_MOSTLY, + hipMemAdviseUnsetReadMostly), + std::make_pair(UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION, + hipMemAdviseSetPreferredLocation), + std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION, + hipMemAdviseUnsetPreferredLocation), + std::make_pair(UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE, + hipMemAdviseSetAccessedBy), + std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE, + hipMemAdviseUnsetAccessedBy), }; - for (auto &FlagPair : URToHIPMemAdviseDeviceFlagsMap) { + for (auto &FlagPair : URToHIPMemAdviseDeviceFlags) { if (URAdviceFlags & FlagPair.first) { UR_CHECK_ERROR(hipMemAdvise(DevPtr, Size, FlagPair.second, Device)); } } - static std::unordered_map - URToHIPMemAdviseHostFlagsMap = { - {UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION_HOST, - hipMemAdviseSetPreferredLocation}, - {UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST, - hipMemAdviseUnsetPreferredLocation}, - {UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_HOST, hipMemAdviseSetAccessedBy}, - {UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_HOST, - hipMemAdviseUnsetAccessedBy}, - }; - - for (auto &FlagPair : URToHIPMemAdviseHostFlagsMap) { + static constexpr std::array URToHIPMemAdviseHostFlags{ + std::make_pair(UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION_HOST, + hipMemAdviseSetPreferredLocation), + std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION_HOST, + hipMemAdviseUnsetPreferredLocation), + std::make_pair(UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_HOST, + hipMemAdviseSetAccessedBy), + std::make_pair(UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_HOST, + hipMemAdviseUnsetAccessedBy), + }; + + for (auto &FlagPair : URToHIPMemAdviseHostFlags) { if (URAdviceFlags & FlagPair.first) { UR_CHECK_ERROR( hipMemAdvise(DevPtr, Size, FlagPair.second, hipCpuDeviceId)); } } - static constexpr std::array UnmappedMemAdviceFlags = - {UR_USM_ADVICE_FLAG_SET_NON_ATOMIC_MOSTLY, - UR_USM_ADVICE_FLAG_CLEAR_NON_ATOMIC_MOSTLY, - UR_USM_ADVICE_FLAG_BIAS_CACHED, UR_USM_ADVICE_FLAG_BIAS_UNCACHED}; - - for (auto &UnmappedFlag : UnmappedMemAdviceFlags) { - if (URAdviceFlags & UnmappedFlag) { - return UR_RESULT_ERROR_INVALID_ENUMERATION; - } + // Handle unmapped memory advice flags + if (URAdviceFlags & + (UR_USM_ADVICE_FLAG_SET_NON_ATOMIC_MOSTLY | + UR_USM_ADVICE_FLAG_CLEAR_NON_ATOMIC_MOSTLY | + UR_USM_ADVICE_FLAG_BIAS_CACHED | UR_USM_ADVICE_FLAG_BIAS_UNCACHED)) { + return UR_RESULT_ERROR_INVALID_ENUMERATION; } return UR_RESULT_SUCCESS; @@ -1390,17 +1391,17 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, #if HIP_VERSION_MAJOR >= 5 UR_ASSERT(pMem && size > 0, UR_RESULT_ERROR_INVALID_VALUE); void *HIPDevicePtr = const_cast(pMem); - - // Passing MEM_ADVISE_SET/MEM_ADVISE_CLEAR_PREFERRED_LOCATION and - // to hipMemAdvise on a GPU device requires the GPU device to report a - // non-zero value for hipDeviceAttributeConcurrentManagedAccess. Therfore, - // ignore memory advise if concurrent managed memory access is not available. - if ((advice & UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION) || - (advice & UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION) || - (advice & UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE) || - (advice & UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE) || - (advice & UR_USM_ADVICE_FLAG_DEFAULT)) { - ur_device_handle_t Device = hQueue->getContext()->getDevice(); + ur_device_handle_t Device = hQueue->getContext()->getDevice(); + + // Passing MEM_ADVISE_SET/MEM_ADVISE_CLEAR_PREFERRED_LOCATION to hipMemAdvise + // on a GPU device requires the GPU device to report a non-zero value for + // hipDeviceAttributeConcurrentManagedAccess. Therefore, ignore the mem advice + // if concurrent managed memory access is not available. + if (advice & (UR_USM_ADVICE_FLAG_SET_PREFERRED_LOCATION | + UR_USM_ADVICE_FLAG_CLEAR_PREFERRED_LOCATION | + UR_USM_ADVICE_FLAG_SET_ACCESSED_BY_DEVICE | + UR_USM_ADVICE_FLAG_CLEAR_ACCESSED_BY_DEVICE | + UR_USM_ADVICE_FLAG_DEFAULT)) { if (!getAttribute(Device, hipDeviceAttributeConcurrentManagedAccess)) { setErrorMessage("mem_advise ignored as device does not support " "concurrent managed access", @@ -1408,11 +1409,13 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, return UR_RESULT_ERROR_ADAPTER_SPECIFIC; } - // TODO: If pMem points to valid system-allocated pageable memory, we should + // If pMem points to valid system-allocated pageable memory, we should // check that the device also has the hipDeviceAttributePageableMemoryAccess // property. } + // NOTE: The hipPointerGetAttribute API is marked as beta, meaning, while this + // is feature complete, it is still open to changes and outstanding issues. unsigned int PointerRangeSize = 0; UR_CHECK_ERROR(hipPointerGetAttribute( &PointerRangeSize, HIP_POINTER_ATTRIBUTE_RANGE_SIZE, @@ -1423,7 +1426,7 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, std::unique_ptr EventPtr{nullptr}; try { - ScopedContext Active(hQueue->getDevice()); + ScopedContext Active(Device); if (phEvent) { EventPtr = @@ -1432,17 +1435,16 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, EventPtr->start(); } + const auto DeviceID = Device->get(); if (advice & UR_USM_ADVICE_FLAG_DEFAULT) { - UR_CHECK_ERROR(hipMemAdvise(pMem, size, hipMemAdviseUnsetReadMostly, - hQueue->getContext()->getDevice()->get())); - UR_CHECK_ERROR(hipMemAdvise(pMem, size, - hipMemAdviseUnsetPreferredLocation, - hQueue->getContext()->getDevice()->get())); - UR_CHECK_ERROR(hipMemAdvise(pMem, size, hipMemAdviseUnsetAccessedBy, - hQueue->getContext()->getDevice()->get())); + UR_CHECK_ERROR( + hipMemAdvise(pMem, size, hipMemAdviseUnsetReadMostly, DeviceID)); + UR_CHECK_ERROR(hipMemAdvise( + pMem, size, hipMemAdviseUnsetPreferredLocation, DeviceID)); + UR_CHECK_ERROR( + hipMemAdvise(pMem, size, hipMemAdviseUnsetAccessedBy, DeviceID)); } else { - Result = setHipMemAdvise(HIPDevicePtr, size, advice, - hQueue->getContext()->getDevice()->get()); + Result = setHipMemAdvise(HIPDevicePtr, size, advice, DeviceID); // UR_RESULT_ERROR_INVALID_ENUMERATION is returned when using a valid but // currently unmapped advice arguments as not supported by this platform. // Therefore, warn the user instead of throwing and aborting the runtime. From 74085833ab3ce28c549a4e5f1d179aec3b927b53 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Wed, 9 Aug 2023 13:05:52 +0100 Subject: [PATCH 04/17] Add HIP device attribute check for managed memory support before calling hipMemAdvise. --- .../unified_runtime/ur/adapters/hip/enqueue.cpp | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp index 00c8e63294ed6..e022b07c9ced7 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp @@ -1393,7 +1393,16 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, void *HIPDevicePtr = const_cast(pMem); ur_device_handle_t Device = hQueue->getContext()->getDevice(); - // Passing MEM_ADVISE_SET/MEM_ADVISE_CLEAR_PREFERRED_LOCATION to hipMemAdvise + // If the device does not support managed memory access, we can't set + // mem_advise. + if (!getAttribute(Device, hipDeviceAttributeManagedMemory)) { + setErrorMessage("mem_advise ignored as device does not support " + " managed memory access", + UR_RESULT_SUCCESS); + return UR_RESULT_ERROR_ADAPTER_SPECIFIC; + } + + // Passing MEM_ADVICE_SET/MEM_ADVICE_CLEAR_PREFERRED_LOCATION to hipMemAdvise // on a GPU device requires the GPU device to report a non-zero value for // hipDeviceAttributeConcurrentManagedAccess. Therefore, ignore the mem advice // if concurrent managed memory access is not available. @@ -1409,9 +1418,10 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, return UR_RESULT_ERROR_ADAPTER_SPECIFIC; } - // If pMem points to valid system-allocated pageable memory, we should + // TODO: If pMem points to valid system-allocated pageable memory, we should // check that the device also has the hipDeviceAttributePageableMemoryAccess - // property. + // property, so that a valid read-only copy can be created on the device. + // This also applies for UR_USM_MEM_ADVICE_SET/MEM_ADVICE_CLEAR_READ_MOSTLY. } // NOTE: The hipPointerGetAttribute API is marked as beta, meaning, while this From 5699baf21f31d8533e9ca937163e775f4eed6006 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Mon, 4 Sep 2023 14:19:44 +0100 Subject: [PATCH 05/17] Feedback - Add _pi_mem_advice aliases for HIP --- sycl/include/sycl/detail/pi.h | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 4ad418b5ccbb7..2abdd67cf81ae 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -145,6 +145,7 @@ // - piextSignalExternalSemaphore // 14.37 Added piextUSMImportExternalPointer and piextUSMReleaseImportedPointer. // 14.38 Change PI_MEM_ADVICE_* values to flags for use in bitwise operations. +// 14.39 Add HIP _pi_mem_advice alises to match the PI_MEM_ADVICE_CUDA* ones. #define _PI_H_VERSION_MAJOR 14 #define _PI_H_VERSION_MINOR 38 @@ -577,6 +578,18 @@ typedef enum { PI_MEM_ADVICE_UNKNOWN = 0x7FFFFFFF, } _pi_mem_advice; +// HIP _pi_mem_advice aliases +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_READ_MOSTLY = PI_MEM_ADVICE_CUDA_SET_READ_MOSTLY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_READ_MOSTLY = PI_MEM_ADVICE_CUDA_UNSET_READ_MOSTLY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION = PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION = PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_ACCESSED_BY = PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY = PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION_HOST = PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION_HOST; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION_HOST = PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_ACCESSED_BY_HOST = PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY_HOST = PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST; + typedef enum { PI_IMAGE_CHANNEL_ORDER_A = 0x10B1, PI_IMAGE_CHANNEL_ORDER_R = 0x10B0, From 4e3a44bcd89c0738258e2f1fbfe40d04c86b6e87 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 5 Sep 2023 16:27:01 +0100 Subject: [PATCH 06/17] Feedback - Rename the memadvise test filename as it applies to more backends --- sycl/test-e2e/USM/{memadvise_cuda.cpp => memadvise_flags.cpp} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename sycl/test-e2e/USM/{memadvise_cuda.cpp => memadvise_flags.cpp} (100%) diff --git a/sycl/test-e2e/USM/memadvise_cuda.cpp b/sycl/test-e2e/USM/memadvise_flags.cpp similarity index 100% rename from sycl/test-e2e/USM/memadvise_cuda.cpp rename to sycl/test-e2e/USM/memadvise_flags.cpp From dc5695e2ee5717d296e508a755c5428544e1064d Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 5 Sep 2023 16:36:43 +0100 Subject: [PATCH 07/17] Feedback - Update the memadvise flags test to include both the CUDA and HIP --- sycl/test-e2e/USM/memadvise_flags.cpp | 48 ++++++++++++++++++--------- 1 file changed, 33 insertions(+), 15 deletions(-) diff --git a/sycl/test-e2e/USM/memadvise_flags.cpp b/sycl/test-e2e/USM/memadvise_flags.cpp index 629cc258b3c22..d6723528877cd 100755 --- a/sycl/test-e2e/USM/memadvise_flags.cpp +++ b/sycl/test-e2e/USM/memadvise_flags.cpp @@ -32,23 +32,41 @@ int main() { return -1; } - // NOTE: PI_MEM_ADVICE_CUDA_* advice values are mapped to the HIP backend too. - std::vector valid_advices{ - PI_MEM_ADVICE_CUDA_SET_READ_MOSTLY, - PI_MEM_ADVICE_CUDA_UNSET_READ_MOSTLY, - PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION, - PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION, - PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY, - PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY, - PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION_HOST, - PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST, - PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST, - PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST, - }; - for (int advice : valid_advices) { - q.mem_advise(ptr, size, advice); + bool isCuda = dev.get_backend() == sycl::backend::ext_oneapi_cuda; + bool isHip = dev.get_backend() == sycl::backend::ext_oneapi_hip; + + std::vector valid_advices; + if (isCuda) { + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_READ_MOSTLY); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_READ_MOSTLY); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION_HOST); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST); + valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST); + } else if (isHip) { + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_SET_READ_MOSTLY); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_UNSET_READ_MOSTLY); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_SET_ACCESSED_BY); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION_HOST); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION_HOST); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_SET_ACCESSED_BY_HOST); + valid_advices.emplace_back(PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY_HOST); + } else { + // Skip + return 0; } + for (int advice : valid_advices) { + q.mem_advise(ptr, size, advice); + } + q.wait_and_throw(); std::cout << "Test passed." << std::endl; return 0; From 499fc4d3476c15668a0047467004b29982feb272 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Mon, 18 Sep 2023 09:58:37 +0100 Subject: [PATCH 08/17] Make mem_advise warn instead of throw for unsupported advice arguments. --- sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp index e022b07c9ced7..e8e614c41d946 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/hip/enqueue.cpp @@ -1461,7 +1461,7 @@ urEnqueueUSMAdvise(ur_queue_handle_t hQueue, const void *pMem, size_t size, if (Result == UR_RESULT_ERROR_INVALID_ENUMERATION) { setErrorMessage("mem_advise is ignored as the advice argument is not " " supported by this device.", - Result); + UR_RESULT_SUCCESS); return UR_RESULT_ERROR_ADAPTER_SPECIFIC; } } From b41ed80c5a3cbcfaa4aabffe1f529b37995a1485 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Mon, 18 Sep 2023 10:22:10 +0100 Subject: [PATCH 09/17] Apply missed clang-format. --- sycl/include/sycl/detail/pi.h | 31 ++++++++++++++++++--------- sycl/test-e2e/USM/memadvise_flags.cpp | 7 +++--- 2 files changed, 25 insertions(+), 13 deletions(-) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index a409bd53dc59c..ffb2ae6ebae6b 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -579,16 +579,27 @@ typedef enum { } _pi_mem_advice; // HIP _pi_mem_advice aliases -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_READ_MOSTLY = PI_MEM_ADVICE_CUDA_SET_READ_MOSTLY; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_READ_MOSTLY = PI_MEM_ADVICE_CUDA_UNSET_READ_MOSTLY; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION = PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION = PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_ACCESSED_BY = PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY = PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION_HOST = PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION_HOST; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION_HOST = PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_ACCESSED_BY_HOST = PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST; -static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY_HOST = PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_READ_MOSTLY = + PI_MEM_ADVICE_CUDA_SET_READ_MOSTLY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_READ_MOSTLY = + PI_MEM_ADVICE_CUDA_UNSET_READ_MOSTLY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION = + PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION = + PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_ACCESSED_BY = + PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY = + PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_PREFERRED_LOCATION_HOST = + PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION_HOST; +static constexpr _pi_mem_advice + PI_MEM_ADVICE_HIP_UNSET_PREFERRED_LOCATION_HOST = + PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_SET_ACCESSED_BY_HOST = + PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST; +static constexpr _pi_mem_advice PI_MEM_ADVICE_HIP_UNSET_ACCESSED_BY_HOST = + PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST; typedef enum { PI_IMAGE_CHANNEL_ORDER_A = 0x10B1, diff --git a/sycl/test-e2e/USM/memadvise_flags.cpp b/sycl/test-e2e/USM/memadvise_flags.cpp index d6723528877cd..b2c0bc8830e33 100755 --- a/sycl/test-e2e/USM/memadvise_flags.cpp +++ b/sycl/test-e2e/USM/memadvise_flags.cpp @@ -44,7 +44,8 @@ int main() { valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY); valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY); valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_PREFERRED_LOCATION_HOST); - valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST); + valid_advices.emplace_back( + PI_MEM_ADVICE_CUDA_UNSET_PREFERRED_LOCATION_HOST); valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_SET_ACCESSED_BY_HOST); valid_advices.emplace_back(PI_MEM_ADVICE_CUDA_UNSET_ACCESSED_BY_HOST); } else if (isHip) { @@ -64,8 +65,8 @@ int main() { } for (int advice : valid_advices) { - q.mem_advise(ptr, size, advice); - } + q.mem_advise(ptr, size, advice); + } q.wait_and_throw(); std::cout << "Test passed." << std::endl; From 0d7516302b2704b9cc0268fbf1cfe923ac9f34fc Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 19 Sep 2023 10:22:14 +0100 Subject: [PATCH 10/17] Update the PI minor to 39 from 38 --- sycl/include/sycl/detail/pi.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index ffb2ae6ebae6b..d14d68f866ff3 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -148,7 +148,7 @@ // 14.39 Add HIP _pi_mem_advice alises to match the PI_MEM_ADVICE_CUDA* ones. #define _PI_H_VERSION_MAJOR 14 -#define _PI_H_VERSION_MINOR 38 +#define _PI_H_VERSION_MINOR 39 #define _PI_STRING_HELPER(a) #a #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b) From a651d6614a476c3a7e76fcaf3deee89ffb483544 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Thu, 21 Sep 2023 16:47:30 +0100 Subject: [PATCH 11/17] Bump minor again after merge commit changes to PI header --- sycl/include/sycl/detail/pi.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 55bfd38d88633..7f24d6febe623 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -149,7 +149,7 @@ // 14.40 Add HIP _pi_mem_advice alises to match the PI_MEM_ADVICE_CUDA* ones. #define _PI_H_VERSION_MAJOR 14 -#define _PI_H_VERSION_MINOR 39 +#define _PI_H_VERSION_MINOR 40 #define _PI_STRING_HELPER(a) #a #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b) From cb6f7b29b9ed7d841e36fb852522f8d800ee4f37 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Thu, 21 Sep 2023 16:52:30 +0100 Subject: [PATCH 12/17] Update test head-comment description. --- sycl/test-e2e/USM/memadvise_flags.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/USM/memadvise_flags.cpp b/sycl/test-e2e/USM/memadvise_flags.cpp index b2c0bc8830e33..542ee7a1ea7a2 100755 --- a/sycl/test-e2e/USM/memadvise_flags.cpp +++ b/sycl/test-e2e/USM/memadvise_flags.cpp @@ -2,7 +2,7 @@ // REQUIRES: cuda || hip_amd // RUN: %{run} %t1.out -//==---------------- memadvise_cuda.cpp ------------------------------------==// +//==---------------- memadvise_flags.cpp -----------------------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From fa7e8cf4f33952ca62fc59d91e0f1d1f5b97e417 Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Wed, 1 Nov 2023 17:51:49 +0000 Subject: [PATCH 13/17] Temporary update CMakeLists to test the UR-HIP adapter changes --- sycl/plugins/unified_runtime/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 690ef378d1cba..1b29a90db68f5 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -4,9 +4,9 @@ # Options to override the default behaviour of the FetchContent to include UR # source code. set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO - "" CACHE STRING "Override the Unified Runtime FetchContent repository") + "https://github.com/GeorgeWeb/unified-runtime.git" CACHE STRING "Override the Unified Runtime FetchContent repository") set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG - "" CACHE STRING "Override the Unified Runtime FetchContent tag") + "be53fb3ba5bf1ac33051456d72c61b6a01d94a72" CACHE STRING "Override the Unified Runtime FetchContent tag") # Options to disable use of FetchContent to include Unified Runtime source code # to improve developer workflow. From 870a7afd21e67e2fe4da22ca60bbec1642ac701b Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 12 Dec 2023 12:44:49 +0000 Subject: [PATCH 14/17] Use updated version of the corresponding UR repo changes --- sycl/plugins/unified_runtime/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 909d23979b30e..8584a2c7706f7 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -6,7 +6,7 @@ set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO "https://github.com/GeorgeWeb/unified-runtime.git" CACHE STRING "Override the Unified Runtime FetchContent repository") set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG - "be53fb3ba5bf1ac33051456d72c61b6a01d94a72" CACHE STRING "Override the Unified Runtime FetchContent tag") + "5a4806ffbf695b6fc89cd5bef337ada8e9b7a573" CACHE STRING "Override the Unified Runtime FetchContent tag") # Options to disable use of FetchContent to include Unified Runtime source code # to improve developer workflow. From 575d5839576bc906252cf75055d8e2ebbb21981f Mon Sep 17 00:00:00 2001 From: "Kenneth Benzie (Benie)" Date: Tue, 9 Jan 2024 10:56:45 +0000 Subject: [PATCH 15/17] [UR] Bump tag to 12a67f56 --- sycl/plugins/unified_runtime/CMakeLists.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index d91991ddfbd27..04a35fd82242f 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -57,13 +57,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT) include(FetchContent) set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit c311fe82256a7bc7f6ddd19cb86c8d555ce401bc - # Merge: eee75a29 d398d4ae + # commit 12a67f56c3c5d08cfac0852d552b4e5fe0452d04 + # Merge: 2b7b827c c10968f5 # Author: Kenneth Benzie (Benie) - # Date: Thu Jan 4 15:12:54 2024 +0000 - # Merge pull request #1222 from sommerlukas/lukas/comgr-include-rocm4 - # [UR][HIP] Fix include for AMD COMGR - set(UNIFIED_RUNTIME_TAG c311fe82256a7bc7f6ddd19cb86c8d555ce401bc) + # Date: Tue Jan 9 10:53:32 2024 +0000 + # Merge pull request #1027 from GeorgeWeb/georgi/hip_memadvise + # [SYCL][HIP] Implement mem_advise for HIP + set(UNIFIED_RUNTIME_TAG 12a67f56c3c5d08cfac0852d552b4e5fe0452d04) if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO) set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}") From 6d2027b07c86b4ffe773ed9a78ff23e7f616435c Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 9 Jan 2024 13:59:57 +0000 Subject: [PATCH 16/17] Fix hip memadvise discard events regression --- sycl/test-e2e/DiscardEvents/discard_events_usm.cpp | 9 +++++++-- .../DiscardEvents/discard_events_usm_ooo_queue.cpp | 13 +++++++------ 2 files changed, 14 insertions(+), 8 deletions(-) diff --git a/sycl/test-e2e/DiscardEvents/discard_events_usm.cpp b/sycl/test-e2e/DiscardEvents/discard_events_usm.cpp index 3b0d86c2d48c8..11288d6620bfd 100644 --- a/sycl/test-e2e/DiscardEvents/discard_events_usm.cpp +++ b/sycl/test-e2e/DiscardEvents/discard_events_usm.cpp @@ -5,6 +5,11 @@ // The test checks that the last parameter is `nullptr` for all PI calls that // should discard events. // {{0|0000000000000000}} is required for various output on Linux and Windows. +// NOTE: piextUSMEnqueuePrefetch and piextUSMEnqueueMemAdvise in the CUDA and +// HIP backends may return a warning result on Windows with error-code +// -996 (PI_ERROR_PLUGIN_SPECIFIC_ERROR) if USM managed memory is not +// supported or if unsupported advice flags are used for the latter API. +// Since it is a warning it is safe to ignore for this test. // // Everything that follows TestQueueOperations() // CHECK: ---> piextUSMEnqueueMemset( @@ -30,7 +35,7 @@ // CHECK-NEXT: pi_event * : {{0|0000000000000000}}[ nullptr ] // // CHECK: ---> piextUSMEnqueueMemAdvise( -// CHECK: ) ---> pi_result : PI_SUCCESS +// CHECK: ) ---> pi_result : {{PI_SUCCESS|-996}} // CHECK-NEXT: [out]pi_event * : {{0|0000000000000000}}[ nullptr ] // // CHECK: ---> piEnqueueKernelLaunch( @@ -75,7 +80,7 @@ // CHECK-NEXT: pi_event * : {{0|0000000000000000}}[ nullptr ] // // CHECK: ---> piextUSMEnqueueMemAdvise( -// CHECK: ) ---> pi_result : PI_SUCCESS +// CHECK: ) ---> pi_result : {{PI_SUCCESS|-996}} // CHECK-NEXT: [out]pi_event * : {{0|0000000000000000}}[ nullptr ] // // CHECK: ---> piEnqueueKernelLaunch( diff --git a/sycl/test-e2e/DiscardEvents/discard_events_usm_ooo_queue.cpp b/sycl/test-e2e/DiscardEvents/discard_events_usm_ooo_queue.cpp index 89f020be5d83b..cfe72db0c1232 100644 --- a/sycl/test-e2e/DiscardEvents/discard_events_usm_ooo_queue.cpp +++ b/sycl/test-e2e/DiscardEvents/discard_events_usm_ooo_queue.cpp @@ -5,10 +5,11 @@ // The test checks that the last parameter is not `nullptr` for all PI calls // that should discard events. // {{0|0000000000000000}} is required for various output on Linux and Windows. -// NOTE: piextUSMEnqueuePrefetch in the CUDA backend may return a warning -// result on Windows with error-code -996 -// (PI_ERROR_PLUGIN_SPECIFIC_ERROR). Since it is a warning it is safe to -// ignore for this test. +// NOTE: piextUSMEnqueuePrefetch and piextUSMEnqueueMemAdvise in the CUDA and +// HIP backends may return a warning result on Windows with error-code +// -996 (PI_ERROR_PLUGIN_SPECIFIC_ERROR) if USM managed memory is not +// supported or if unsupported advice flags are used for the latter API. +// Since it is a warning it is safe to ignore for this test. // // Everything that follows TestQueueOperations() // CHECK: ---> piextUSMEnqueueMemset( @@ -40,7 +41,7 @@ // // CHECK: ---> piextUSMEnqueueMemAdvise( // CHECK-NOT: pi_event * : {{0|0000000000000000}}[ nullptr ] -// CHECK: ---> pi_result : PI_SUCCESS +// CHECK: ---> pi_result : {{PI_SUCCESS|-996}} // // CHECK: ---> piEnqueueEventsWaitWithBarrier( // CHECK-NOT: pi_event * : {{0|0000000000000000}}[ nullptr ] @@ -97,7 +98,7 @@ // // CHECK: ---> piextUSMEnqueueMemAdvise( // CHECK-NOT: pi_event * : {{0|0000000000000000}}[ nullptr ] -// CHECK: ---> pi_result : PI_SUCCESS +// CHECK: ---> pi_result : {{PI_SUCCESS|-996}} // // CHECK: ---> piEnqueueEventsWaitWithBarrier( // CHECK-NOT: pi_event * : {{0|0000000000000000}}[ nullptr ] From 56abbdb097c78958782bc99f65c88892b2645b8f Mon Sep 17 00:00:00 2001 From: Georgi Mirazchiyski Date: Tue, 9 Jan 2024 16:07:21 +0000 Subject: [PATCH 17/17] Unset the UR tag override cmake variable --- sycl/plugins/unified_runtime/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 04a35fd82242f..46bd16b592553 100644 --- a/sycl/plugins/unified_runtime/CMakeLists.txt +++ b/sycl/plugins/unified_runtime/CMakeLists.txt @@ -4,9 +4,9 @@ # Options to override the default behaviour of the FetchContent to include UR # source code. set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO - "https://github.com/GeorgeWeb/unified-runtime.git" CACHE STRING "Override the Unified Runtime FetchContent repository") + "" CACHE STRING "Override the Unified Runtime FetchContent repository") set(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_TAG - "5a4806ffbf695b6fc89cd5bef337ada8e9b7a573" CACHE STRING "Override the Unified Runtime FetchContent tag") + "" CACHE STRING "Override the Unified Runtime FetchContent tag") # Options to disable use of FetchContent to include Unified Runtime source code # to improve developer workflow.