From deeefffff5c8de58a172faef1e50bb42b7214492 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Thu, 7 Mar 2019 14:58:49 -0500
Subject: [PATCH] P4 to Git Change 1752994 by kjayapra@2_HIPWS_SL_D32 on
 2019/03/07 14:06:54

	SWDEV-145570 - Implementation of ihipMemsetD32 and ihipMemsetD32Async
	       HIP-HCC Info: https://github.com/ROCm-Developer-Tools/HIP/pull/933

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#11 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#12 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#44 edit


[ROCm/clr commit: 1f45d27b9f1d94a5f0142c28b180ce3f7adf6b4c]
---
 projects/clr/hipamd/api/hip/hip_hcc.def.in |  2 +
 projects/clr/hipamd/api/hip/hip_hcc.map.in |  2 +
 projects/clr/hipamd/api/hip/hip_memory.cpp | 45 ++++++++++++----------
 3 files changed, 29 insertions(+), 20 deletions(-)

diff --git a/projects/clr/hipamd/api/hip/hip_hcc.def.in b/projects/clr/hipamd/api/hip/hip_hcc.def.in
index 576f840695..f0b4560fd6 100644
--- a/projects/clr/hipamd/api/hip/hip_hcc.def.in
+++ b/projects/clr/hipamd/api/hip/hip_hcc.def.in
@@ -95,8 +95,10 @@ hipMemGetAddressRange
 hipMemGetInfo
 hipMemPtrGetInfo
 hipMemset
+hipMemsetD32
 hipMemset2D
 hipMemsetAsync
+hipMemsetD32Async
 hipMemset2DAsync
 hipMemsetD8
 hipMemset3D
diff --git a/projects/clr/hipamd/api/hip/hip_hcc.map.in b/projects/clr/hipamd/api/hip/hip_hcc.map.in
index 8f033617be..b6d16035fa 100644
--- a/projects/clr/hipamd/api/hip/hip_hcc.map.in
+++ b/projects/clr/hipamd/api/hip/hip_hcc.map.in
@@ -96,8 +96,10 @@ global:
     hipMemGetInfo;
     hipMemPtrGetInfo;
     hipMemset;
+    hipMemsetD32;
     hipMemset2D;
     hipMemsetAsync;
+    hipMemsetD32Async;
     hipMemset2DAsync;
     hipMemsetD8;
     hipMemset3D;
diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp
index 587bbbc96e..2a73d98378 100644
--- a/projects/clr/hipamd/api/hip/hip_memory.cpp
+++ b/projects/clr/hipamd/api/hip/hip_memory.cpp
@@ -112,24 +112,33 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   return hipSuccess;
 }
 
-hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, amd::HostQueue& queue,
-                      bool isAsync = false) {
+hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes,
+                      hipStream_t stream, bool isAsync = false) {
 
   if (dst == nullptr) {
     return hipErrorInvalidValue;
   }
 
   size_t offset = 0;
+  amd::HostQueue* queue = nullptr;
   amd::Memory* memory = getMemoryObject(dst, offset);
 
+  if (stream == nullptr) {
+    hip::syncStreams();
+    queue = hip::getNullStream();
+  } else {
+    hip::getNullStream()->finish();
+    queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+  }
+
   if (memory != nullptr) {
     // Device memory
     amd::Command::EventWaitList waitList;
     amd::Coord3D fillOffset(offset, 0, 0);
     amd::Coord3D fillSize(sizeBytes, 1, 1);
     amd::FillMemoryCommand* command =
-        new amd::FillMemoryCommand(queue, CL_COMMAND_FILL_BUFFER, waitList, *memory->asBuffer(),
-                                   &value, sizeof(char), fillOffset, fillSize);
+        new amd::FillMemoryCommand(*queue, CL_COMMAND_FILL_BUFFER, waitList, *memory->asBuffer(),
+                                   &value, valueSize, fillOffset, fillSize);
 
     if (command == nullptr) {
       return hipErrorOutOfMemory;
@@ -196,26 +205,25 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
   HIP_INIT_API(dst, value, sizeBytes, stream);
 
-  amd::HostQueue* queue;
+  HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, stream, true));
+}
 
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
-  }
+hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) {
+  HIP_INIT_API(dst, value, count, stream);
 
-  HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue, true));
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), stream, true));
 }
 
 hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
   HIP_INIT_API(dst, value, sizeBytes);
 
-  hip::syncStreams();
-  amd::HostQueue* queue = hip::getNullStream();
+  HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
+}
 
-  HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue));
+hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) {
+  HIP_INIT_API(dst, value, count);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), nullptr));
 }
 
 hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) {
@@ -360,10 +368,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent)
   void *dst = pitchedDevPtr.ptr;
   size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
 
-  hip::syncStreams();
-  amd::HostQueue* queue = hip::getNullStream();
-
-  HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue));
+  HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
 }
 
 hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {