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: 1f45d27b9f]

projects/clr/hipamd/api/hip/hip_hcc.def.in

@@ -95,8 +95,10 @@ hipMemGetAddressRange
 hipMemGetInfo
 hipMemPtrGetInfo
 hipMemset
+hipMemsetD32
 hipMemset2D
 hipMemsetAsync
+hipMemsetD32Async
 hipMemset2DAsync
 hipMemsetD8
 hipMemset3D

projects/clr/hipamd/api/hip/hip_hcc.map.in

@@ -96,8 +96,10 @@ global:
     hipMemGetInfo;
     hipMemPtrGetInfo;
     hipMemset;
+    hipMemsetD32;
     hipMemset2D;
     hipMemsetAsync;
+    hipMemsetD32Async;
     hipMemset2DAsync;
     hipMemsetD8;
     hipMemset3D;

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,
+hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes,
-                      bool isAsync = false) {
+                      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(),
+        new amd::FillMemoryCommand(*queue, CL_COMMAND_FILL_BUFFER, waitList, *memory->asBuffer(),
-                                   &value, sizeof(char), fillOffset, fillSize);
+                                   &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) {
+hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count, hipStream_t stream) {
-    hip::syncStreams();
+  HIP_INIT_API(dst, value, count, stream);
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
-  }
 
-  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();
+  HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
-  amd::HostQueue* queue = hip::getNullStream();
+}
 
-  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();
+  HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
-  amd::HostQueue* queue = hip::getNullStream();
-
-  HIP_RETURN(ihipMemset(dst, value, sizeBytes, *queue));
 }
 
 hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {