From 9d76c76746a5fdbae438de3d5a2e11d169aaa7fc Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Fri, 15 May 2020 18:08:26 -0400 Subject: [PATCH] Add support for hipExtStreamCreateWithCUMask API Change-Id: I369d0eaca493821c4badc6b18ac02daa2fddc95f [ROCm/clr commit: 3a30b9eb8dc75502fd807bdac80cb422ec1c0288] --- .../include/hip/hcc_detail/hip_runtime_api.h | 22 ++++++++++++++ projects/clr/hipamd/rocclr/hip_hcc.def.in | 1 + projects/clr/hipamd/rocclr/hip_hcc.map.in | 1 + projects/clr/hipamd/rocclr/hip_internal.hpp | 4 ++- projects/clr/hipamd/rocclr/hip_stream.cpp | 29 +++++++++++++++---- 5 files changed, 51 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 744885923f..d908f72a04 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -849,6 +849,28 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags); hipError_t hipStreamGetPriority(hipStream_t stream, int* priority); +/** + * @brief Create an asynchronous stream with the specified CU mask. + * + * @param[in, out] stream Pointer to new stream + * @param[in ] cuMaskSize Size of CU mask bit array passed in. + * @param[in ] cuMask Bit-vector representing the CU mask. Each active bit represents using one CU. + * The first 32 bits represent the first 32 CUs, and so on. If its size is greater than physical + * CU number (i.e., multiProcessorCount member of hipDeviceProp_t), the extra elements are ignored. + * It is user's responsibility to make sure the input is meaningful. + * @return #hipSuccess, #hipErrorInvalidHandle, #hipErrorInvalidValue + * + * Create a new asynchronous stream with the specified CU mask. @p stream returns an opaque handle + * that can be used to reference the newly created stream in subsequent hipStream* commands. The + * stream is allocated on the heap and will remain allocated even if the handle goes out-of-scope. + * To release the memory used by the stream, application must call hipStreamDestroy. + * + * + * @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy + */ +hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, const uint32_t* cuMask); + + /** * Stream CallBack struct */ diff --git a/projects/clr/hipamd/rocclr/hip_hcc.def.in b/projects/clr/hipamd/rocclr/hip_hcc.def.in index 579608e685..53bdfb3065 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.def.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.def.in @@ -252,3 +252,4 @@ hipTexObjectDestroy hipTexObjectGetResourceDesc hipTexObjectGetResourceViewDesc hipTexObjectGetTextureDesc +hipExtStreamCreateWithCUMask diff --git a/projects/clr/hipamd/rocclr/hip_hcc.map.in b/projects/clr/hipamd/rocclr/hip_hcc.map.in index 19da8a6991..02a704bf88 100755 --- a/projects/clr/hipamd/rocclr/hip_hcc.map.in +++ b/projects/clr/hipamd/rocclr/hip_hcc.map.in @@ -258,6 +258,7 @@ global: hipInitActivityCallback*; hipEnableActivityCallback*; hipGetCmdName*; + hipExtStreamCreateWithCUMask; }; local: *; diff --git a/projects/clr/hipamd/rocclr/hip_internal.hpp b/projects/clr/hipamd/rocclr/hip_internal.hpp index 4cc0dadd8a..492842c32b 100755 --- a/projects/clr/hipamd/rocclr/hip_internal.hpp +++ b/projects/clr/hipamd/rocclr/hip_internal.hpp @@ -87,9 +87,11 @@ namespace hip { amd::CommandQueue::Priority priority_; unsigned int flags_; bool null_; + const std::vector cuMask_; public: - Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false); + Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false, + const std::vector& cuMask = {}); /// Creates the hip stream object, including AMD host queue bool Create(); diff --git a/projects/clr/hipamd/rocclr/hip_stream.cpp b/projects/clr/hipamd/rocclr/hip_stream.cpp index 3bd7d343f7..23b89e71ca 100644 --- a/projects/clr/hipamd/rocclr/hip_stream.cpp +++ b/projects/clr/hipamd/rocclr/hip_stream.cpp @@ -44,15 +44,15 @@ namespace hip { // ================================================================================================ Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p, - unsigned int f, bool null_stream) + unsigned int f, bool null_stream, const std::vector& cuMask) : queue_(nullptr), lock_("Stream Callback lock"), device_(dev), - priority_(p), flags_(f), null_(null_stream) {} + priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {} // ================================================================================================ bool Stream::Create() { cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE; queue_ = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties, - amd::CommandQueue::RealTimeDisabled, priority_); + amd::CommandQueue::RealTimeDisabled, priority_, cuMask_); // Create a host queue bool result = (queue_ != nullptr) ? queue_->create() : false; // Insert just created stream into the list of the blocking queues @@ -160,8 +160,9 @@ void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, // ================================================================================================ static hipError_t ihipStreamCreate(hipStream_t* stream, - unsigned int flags, amd::CommandQueue::Priority priority) { - hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags); + unsigned int flags, amd::CommandQueue::Priority priority, + const std::vector& cuMask = {}) { + hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask); if (hStream == nullptr) { return hipErrorOutOfMemory; @@ -311,3 +312,21 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback HIP_RETURN(hipSuccess); } + +// ================================================================================================ +hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, + const uint32_t* cuMask) { + HIP_INIT_API(hipExtStreamCreateWithCUMask, stream, cuMaskSize, cuMask); + + if (stream == nullptr) { + HIP_RETURN(hipErrorInvalidHandle); + } + if (cuMaskSize == 0 || cuMask == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + const std::vector cuMaskv(cuMask, cuMask + cuMaskSize); + + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv)); +} +