diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 171bab780e1a3..010c59dd3c9d6 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -146,9 +146,10 @@ // 14.37 Added piextUSMImportExternalPointer and piextUSMReleaseImportedPointer. // 14.38 Change PI_MEM_ADVICE_* values to flags for use in bitwise operations. // 14.39 Added PI_EXT_INTEL_DEVICE_INFO_ESIMD_SUPPORT device info query. +// 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) @@ -580,6 +581,29 @@ 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, diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt index 7eaca229b619c..46bd16b592553 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 2d42d77d5b931b4558842bad32cf9dad80e17b3c - # Merge: 9f14fedb 2a960baf + # commit 12a67f56c3c5d08cfac0852d552b4e5fe0452d04 + # Merge: 2b7b827c c10968f5 # Author: Kenneth Benzie (Benie) - # Date: Mon Jan 8 13:40:40 2024 +0000 - # Merge pull request #1229 from oneapi-src/revert-1228-revert-984-ext_oneapi_queue_priority-hip - # Revert "Revert "[HIP] Implement ext_oneapi_queue_priority"" - set(UNIFIED_RUNTIME_TAG 2d42d77d5b931b4558842bad32cf9dad80e17b3c) + # 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}") diff --git a/sycl/plugins/unified_runtime/pi2ur.hpp b/sycl/plugins/unified_runtime/pi2ur.hpp index 6e36b21c31b89..d3051c47bd93b 100644 --- a/sycl/plugins/unified_runtime/pi2ur.hpp +++ b/sycl/plugins/unified_runtime/pi2ur.hpp @@ -3358,6 +3358,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/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 ] diff --git a/sycl/test-e2e/USM/memadvise_cuda.cpp b/sycl/test-e2e/USM/memadvise_cuda.cpp deleted file mode 100644 index cacd5be1ae7bf..0000000000000 --- a/sycl/test-e2e/USM/memadvise_cuda.cpp +++ /dev/null @@ -1,54 +0,0 @@ -// RUN: %{build} -o %t1.out -// REQUIRES: cuda -// RUN: %{run} %t1.out - -//==---------------- memadvise_cuda.cpp ------------------------------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include -#include -#include - -using namespace sycl; - -int main() { - const size_t size = 100; - queue q; - auto dev = q.get_device(); - auto ctx = q.get_context(); - if (!dev.get_info()) { - std::cout << "Shared USM is not supported. Skipping test." << std::endl; - return 0; - } - - void *ptr = malloc_shared(size, dev, ctx); - if (ptr == nullptr) { - std::cout << "Allocation failed!" << std::endl; - return -1; - } - - 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); - } - - q.wait_and_throw(); - std::cout << "Test passed." << std::endl; - return 0; -} diff --git a/sycl/test-e2e/USM/memadvise_flags.cpp b/sycl/test-e2e/USM/memadvise_flags.cpp new file mode 100755 index 0000000000000..542ee7a1ea7a2 --- /dev/null +++ b/sycl/test-e2e/USM/memadvise_flags.cpp @@ -0,0 +1,74 @@ +// RUN: %{build} -o %t1.out +// REQUIRES: cuda || hip_amd +// RUN: %{run} %t1.out + +//==---------------- 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. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include + +using namespace sycl; + +int main() { + const size_t size = 100; + queue q; + auto dev = q.get_device(); + auto ctx = q.get_context(); + if (!dev.get_info()) { + std::cout << "Shared USM is not supported. Skipping test." << std::endl; + return 0; + } + + void *ptr = malloc_shared(size, dev, ctx); + if (ptr == nullptr) { + std::cout << "Allocation failed!" << std::endl; + return -1; + } + + 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; +}