Add support for hipExtStreamCreateWithCUMask API

Change-Id: I369d0eaca493821c4badc6b18ac02daa2fddc95f


[ROCm/clr commit: 3a30b9eb8d]

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
  */

projects/clr/hipamd/rocclr/hip_hcc.def.in

@@ -252,3 +252,4 @@ hipTexObjectDestroy
 hipTexObjectGetResourceDesc
 hipTexObjectGetResourceViewDesc
 hipTexObjectGetTextureDesc
+hipExtStreamCreateWithCUMask

projects/clr/hipamd/rocclr/hip_hcc.map.in

@@ -258,6 +258,7 @@ global:
     hipInitActivityCallback*;
     hipEnableActivityCallback*;
     hipGetCmdName*;
+    hipExtStreamCreateWithCUMask;
     };
 local:
     *;

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<uint32_t> 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<uint32_t>& cuMask = {});
 
     /// Creates the hip stream object, including AMD host queue
     bool Create();

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<uint32_t>& 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<uint32_t>& 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<uint32_t> cuMaskv(cuMask, cuMask + cuMaskSize);
+
+  HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv));
+}
+