From b99ebf8d006b0e7bb53ed7e15482f952d646a1be Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <Sarbojit.Sarkar@amd.com>
Date: Wed, 16 Jun 2021 10:16:18 +0000
Subject: [PATCH] SWDEV-284435 - StreamOps mapping with cuda drv API

Change-Id: I6ed9196cfe59d9004338206f6dd8d6fb367b3ffa


[ROCm/clr commit: 7e3c17c023f6381a5c6c8133b724d11b6728632b]
---
 .../nvidia_detail/nvidia_hip_runtime_api.h    | 48 +++++++++++++++++++
 1 file changed, 48 insertions(+)

diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
index 66e4743abd..5d1813210e 100644
--- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
@@ -81,6 +81,10 @@ typedef enum hipMemoryAdvise {
 #define HIP_C_32F CUDA_C_32F
 #define HIP_C_64F CUDA_C_64F
 
+// hip stream operation masks
+#define STREAM_OPS_WAIT_MASK_32 0xFFFFFFFF
+#define STREAM_OPS_WAIT_MASK_64 0xFFFFFFFFFFFFFFFF
+
 // hipLibraryPropertyType
 #define hipLibraryPropertyType libraryPropertyType
 #define HIP_LIBRARY_MAJOR_VERSION MAJOR_VERSION
@@ -2190,6 +2194,50 @@ inline static hipError_t hipArray3DCreate(hiparray* pHandle,
    return hipCUResultTohipError(cuArray3DCreate(pHandle, pAllocateArray));
 }
 
+inline static hipError_t hipStreamWriteValue32(hipStream_t stream,
+                                               void* ptr, int32_t value, unsigned int flags) {
+   if (value < 0) {
+     printf("Warning! value is negative, CUDA accept positive values\n");
+   }
+   return hipCUResultTohipError(cuStreamWriteValue32(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                             reinterpret_cast<cuuint32_t>(value),flags));
+}
+
+inline static hipError_t hipStreamWriteValue64(hipStream_t stream,
+                                               void* ptr, int64_t value, unsigned int flags) {
+   if (value < 0) {
+     printf("Warning! value is negative, CUDA accept positive values\n");
+   }
+   return hipCUResultTohipError(cuStreamWriteValue64(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                    reinterpret_cast<cuuint64_t>(value),flags));
+}
+
+inline static hipError_t hipStreamWaitValue32(hipStream_t stream,
+                                              void* ptr, int32_t value, unsigned int flags,
+                                              uint32_t mask)) {
+   if (value < 0) {
+     printf("Warning! value is negative, CUDA accept positive values\n");
+   }
+   if (mask != STREAM_OPS_WAIT_MASK_32) {
+     printf("Warning! mask will not have impact as CUDA ignores it.\n");
+   }
+   return hipCUResultTohipError(cuStreamWaitValue32(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                    reinterpret_cast<cuuint32_t>(value),flags));
+}
+
+inline static hipError_t hipStreamWaitValue64(hipStream_t stream,
+                                              void* ptr, int64_t value, unsigned int flags,
+                                              uint64_t mask)) {
+   if (value < 0) {
+     printf("Warning! value is negative, CUDA accept positive values\n");
+   }
+   if (mask != STREAM_OPS_WAIT_MASK_64) {
+     printf("Warning! mask will not have impact as CUDA ignores it.\n");
+   }
+   return hipCUResultTohipError(cuStreamWaitValue64(stream, reinterpret_cast<CUdeviceptr>(ptr),
+                                                    reinterpret_cast<cuuint64_t>(value),flags));
+}
+
 #endif  //__CUDACC__
 
 #endif  // HIP_INCLUDE_HIP_NVIDIA_DETAIL_HIP_RUNTIME_API_H