From 250ecb0bdb4ff2e40cbf117016c281e68de0bf30 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Tue, 8 Oct 2019 14:55:27 -0400
Subject: [PATCH] P4 to Git Change 2010058 by kjayapra@0_HIPWS_LNX1_ROCM on
 2019/10/08 14:44:53

	SWDEV-144570 - Implementation of APIs hipMemsetD8Async, hipMemAllocPitch, hipMemAllocHost and template for hipMemsetD16, hipMemsetD16Async.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#30 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#29 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#75 edit
---
 hipamd/api/hip/hip_hcc.def.in |  5 ++++
 hipamd/api/hip/hip_hcc.map.in |  5 ++++
 hipamd/api/hip/hip_memory.cpp | 52 +++++++++++++++++++++++++++++++----
 3 files changed, 56 insertions(+), 6 deletions(-)

diff --git a/hipamd/api/hip/hip_hcc.def.in b/hipamd/api/hip/hip_hcc.def.in
index 70eaa66870..45aa1503cc 100644
--- a/hipamd/api/hip/hip_hcc.def.in
+++ b/hipamd/api/hip/hip_hcc.def.in
@@ -60,6 +60,7 @@ hipGetDeviceProperties
 hipGetErrorName
 hipGetErrorString
 hipGetLastError
+hipMemAllocHost
 hipHostAlloc
 hipHostFree
 hipHostGetDevicePointer
@@ -77,6 +78,7 @@ hipMalloc3DArray
 hipArrayCreate
 hipArray3DCreate
 hipMallocArray
+hipMemAllocPitch
 hipMallocPitch
 hipMemcpy
 hipMemcpyParam2D
@@ -105,9 +107,12 @@ hipGetSymbolSize
 hipMemGetInfo
 hipMemPtrGetInfo
 hipMemset
+hipMemsetD16
 hipMemsetD32
 hipMemset2D
 hipMemsetAsync
+hipMemsetD8Async
+hipMemsetD16Async
 hipMemsetD32Async
 hipMemset2DAsync
 hipMemsetD8
diff --git a/hipamd/api/hip/hip_hcc.map.in b/hipamd/api/hip/hip_hcc.map.in
index 217164cafb..8535ab72f1 100644
--- a/hipamd/api/hip/hip_hcc.map.in
+++ b/hipamd/api/hip/hip_hcc.map.in
@@ -61,6 +61,7 @@ global:
     hipGetErrorName;
     hipGetErrorString;
     hipGetLastError;
+    hipMemAllocHost;
     hipHostAlloc;
     hipHostFree;
     hipHostGetDevicePointer;
@@ -79,6 +80,7 @@ global:
     hipArray3DCreate;
     hipMallocArray;
     hipMallocPitch;
+    hipMemAllocPitch;
     hipMemcpy;
     hipMemcpyParam2D;
     hipMemcpy2D;
@@ -106,9 +108,12 @@ global:
     hipMemGetInfo;
     hipMemPtrGetInfo;
     hipMemset;
+    hipMemsetD16;
     hipMemsetD32;
     hipMemset2D;
     hipMemsetAsync;
+    hipMemsetD8Async;
+    hipMemsetD16Async;
     hipMemsetD32Async;
     hipMemset2DAsync;
     hipMemsetD8;
diff --git a/hipamd/api/hip/hip_memory.cpp b/hipamd/api/hip/hip_memory.cpp
index 5abfb7e365..174bd14e5b 100644
--- a/hipamd/api/hip/hip_memory.cpp
+++ b/hipamd/api/hip/hip_memory.cpp
@@ -269,6 +269,14 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
   HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, stream, true));
 }
 
+hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t count, hipStream_t stream){
+  HIP_INIT_API(hipMemsetD16Async, dst, value, count, stream);
+
+  assert(0 && "Unimplemented");
+
+  HIP_RETURN(hipErrorUnknown);
+}
+
 hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) {
   HIP_INIT_API(hipMemsetD32Async, dst, value, count, stream);
 
@@ -281,6 +289,14 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
   HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
 }
 
+hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t count){
+  HIP_INIT_API(hipMemsetD16, dst, value, count);
+
+  assert(0 && "Unimplemented");
+
+  HIP_RETURN(hipErrorUnknown);
+}
+
 hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) {
   HIP_INIT_API(hipMemsetD32, dst, value, count);
 
@@ -1341,16 +1357,14 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
   HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue, true));
 }
 
-hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
-  HIP_INIT_API(hipMemsetD8, dst, value, sizeBytes);
+hipError_t ihipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes, hipStream_t stream) {
 
   if (dst == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+    return hipErrorInvalidValue;
   }
 
-  hip::syncStreams();
-  amd::HostQueue* queue = hip::getNullStream();
   size_t offset = 0;
+  amd::HostQueue* queue = hip::getQueue(stream);
   amd::Command::EventWaitList waitList;
   amd::Memory* memory = getMemoryObject(dst, offset);
   if (memory != nullptr) {
@@ -1373,7 +1387,33 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
     memset(dst, value, sizeBytes);
   }
 
-  HIP_RETURN(hipSuccess);
+  return hipSuccess;
+}
+
+hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
+  HIP_INIT_API(hipMemsetD8, dst, value, sizeBytes);
+
+  HIP_RETURN(ihipMemsetD8(dst, value, sizeBytes, nullptr));
+}
+
+hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes,
+                            hipStream_t stream) {
+  HIP_INIT_API(hipMemsetD8Async, dst, value, sizeBytes, stream);
+
+  HIP_RETURN(ihipMemsetD8(dst, value, sizeBytes, stream));
+}
+
+hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes,
+                            size_t height, unsigned int elementSizeBytes) {
+  HIP_INIT_API(hipMemAllocPitch, dptr, pitch, widthInBytes, height, elementSizeBytes);
+
+  HIP_RETURN(hipMallocPitch(dptr, pitch, widthInBytes, height));
+}
+
+hipError_t hipMemAllocHost(void** ptr, size_t size) {
+  HIP_INIT_API(hipMemAllocHost, ptr, size);
+
+  HIP_RETURN(hipHostMalloc(ptr, size, 0));
 }
 
 hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* dev_ptr) {