P4 to Git Change 2015320 by vsytchen@vsytchen-hip-win10 on 2019/10/17 14:00:39

SWDEV-207100 - [HIP CQE][HIPonPAL][WIN][QR] 5 hiptests failed in 19H1 Windows on all ASICs

	1. Reshuffle locations of the hipMemset functions to make them all next to each other.
	2. Update the declarations of hipMemsetD8, hipMemsetD8Async, hipMemsetD16, hipMemsetD16Async. These functions are type aware and take in as their third argument the number of elements in the buffer, not the buffer size. Change the name of this argument from sizeBytes to count to align with the above description. Changes for the header are tracked here https://github.com/ROCm-Developer-Tools/HIP/pull/1544
	3. Add the actual implementation of hipMemsetD8, hipMemsetD8Async, hipMemsetD16, hipMemsetD16Async.
	4. Remove ihipMemset2D() as it is essentially a copy of ihipMemset(). Change hipMemset2D()/hipMemset2DAsync() to use ihipMemset().
	5. Implement hipMemset3DAsync().
	6. Update the test script to pick up the updated command line options for hipMemset and hipMemset3D.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#32 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#30 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#78 edit
... //depot/stg/opencl/drivers/opencl/make/hip.git/tests/scripts/hip_runtimeapi_tests.txt#13 edit


[ROCm/clr commit: 93ad67f8a6]

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

@@ -107,16 +107,17 @@ hipGetSymbolSize
 hipMemGetInfo
 hipMemPtrGetInfo
 hipMemset
-hipMemsetD16
-hipMemsetD32
-hipMemset2D
 hipMemsetAsync
-hipMemsetD8Async
-hipMemsetD16Async
-hipMemsetD32Async
-hipMemset2DAsync
 hipMemsetD8
+hipMemsetD8Async
+hipMemsetD16
+hipMemsetD16Async
+hipMemsetD32
+hipMemsetD32Async
+hipMemset2D
+hipMemset2DAsync
 hipMemset3D
+hipMemset3DAsync
 hipModuleGetFunction
 hipModuleGetGlobal
 hipModuleGetTexRef

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

@@ -108,16 +108,17 @@ global:
     hipMemGetInfo;
     hipMemPtrGetInfo;
     hipMemset;
-    hipMemsetD16;
-    hipMemsetD32;
-    hipMemset2D;
     hipMemsetAsync;
-    hipMemsetD8Async;
-    hipMemsetD16Async;
-    hipMemsetD32Async;
-    hipMemset2DAsync;
     hipMemsetD8;
+    hipMemsetD8Async;
+    hipMemsetD16;
+    hipMemsetD16Async;
+    hipMemsetD32;
+    hipMemsetD32Async;
+    hipMemset2D;
+    hipMemset2DAsync;
     hipMemset3D;
+    hipMemset3DAsync;
     hipModuleGetFunction;
     hipModuleGetGlobal;
     hipModuleGetTexRef;

projects/clr/hipamd/api/hip/hip_memory.cpp

@@ -160,47 +160,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   return hipSuccess;
 }
 
-hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes,
-                      hipStream_t stream, bool isAsync = false) {
-  if (sizeBytes == 0) {
-    // Skip if nothing needs filling.
-    return hipSuccess;
-  }
-
-  if (dst == nullptr) {
-    return hipErrorInvalidValue;
-  }
-
-  size_t offset = 0;
-  amd::HostQueue* queue = hip::getQueue(stream);
-  amd::Memory* memory = getMemoryObject(dst, offset);
-
-  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, valueSize, fillOffset, fillSize);
-
-    if (command == nullptr) {
-      return hipErrorOutOfMemory;
-    }
-
-    command->enqueue();
-    if (!isAsync) {
-      command->awaitCompletion();
-    }
-    command->release();
-  } else {
-    // Host alloced memory
-    memset(dst, value, sizeBytes);
-  }
-
-  return hipSuccess;
-}
-
 hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) {
   HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags);
 
@@ -265,46 +224,6 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
   HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue));
 }
 
-hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
-  HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream);
-
-  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);
-
-  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(hipMemset, dst, value, 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);
-
-  HIP_RETURN(ihipMemset(dst, value, sizeof(int), count * sizeof(int), nullptr));
-}
-
 hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) {
   HIP_INIT_API(hipMemPtrGetInfo, ptr, size);
 
@@ -441,15 +360,6 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
   HIP_RETURN(status);
 }
 
-hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) {
-  HIP_INIT_API(hipMemset3D, pitchedDevPtr, value, &extent);
-
-  void *dst = pitchedDevPtr.ptr;
-  size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
-
-  HIP_RETURN(ihipMemset(dst, value, sizeof(char), sizeBytes, nullptr));
-}
-
 hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
   HIP_INIT_API(hipArrayCreate, array, pAllocateArray);
 
@@ -1304,16 +1214,19 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
   HIP_RETURN(hipSuccess);
 }
 
-hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height,
-                        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 (sizeBytes == 0) {
+    // Skip if nothing needs filling.
+    return hipSuccess;
+  }
 
   if (dst == nullptr) {
     return hipErrorInvalidValue;
   }
 
-  size_t sizeBytes = pitch * height;
   size_t offset = 0;
-
+  amd::HostQueue* queue = hip::getQueue(stream);
   amd::Memory* memory = getMemoryObject(dst, offset);
 
   if (memory != nullptr) {
@@ -1321,18 +1234,16 @@ hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
     amd::Command::EventWaitList waitList;
     amd::Coord3D fillOffset(offset, 0, 0);
     amd::Coord3D fillSize(sizeBytes, 1, 1);
-
-    // TODO: Byte copies are inefficient. Combine multiple writes inside runtime
     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;
     }
 
     command->enqueue();
-    if(!isAsync) {
+    if (!isAsync) {
       command->awaitCompletion();
     }
     command->release();
@@ -1344,67 +1255,86 @@ hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
   return hipSuccess;
 }
 
+hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
+  HIP_INIT_API(hipMemset, dst, value, sizeBytes);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), sizeBytes, nullptr));
+}
+
+hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
+  HIP_INIT_API(hipMemsetAsync, dst, value, sizeBytes, stream);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), sizeBytes, stream, true));
+}
+
+hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t count) {
+  HIP_INIT_API(hipMemsetD8, dst, value, count);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), count * sizeof(int8_t), nullptr));
+}
+
+hipError_t hipMemsetD8Async(hipDeviceptr_t dst, unsigned char value, size_t count,
+                            hipStream_t stream) {
+  HIP_INIT_API(hipMemsetD8Async, dst, value, count, stream);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), count * sizeof(int8_t), stream, true));
+}
+
+hipError_t hipMemsetD16(hipDeviceptr_t dst, unsigned short value, size_t count) {
+  HIP_INIT_API(hipMemsetD16, dst, value, count);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int16_t), count * sizeof(int16_t), nullptr));
+}
+
+hipError_t hipMemsetD16Async(hipDeviceptr_t dst, unsigned short value, size_t count,
+                             hipStream_t stream) {
+  HIP_INIT_API(hipMemsetD16Async, dst, value, count, stream);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int16_t), count * sizeof(int16_t), stream, true));
+}
+
+hipError_t hipMemsetD32(hipDeviceptr_t dst, int value, size_t count) {
+  HIP_INIT_API(hipMemsetD32, dst, value, count);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int32_t), count * sizeof(int32_t), nullptr));
+}
+
+hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count,
+                             hipStream_t stream) {
+  HIP_INIT_API(hipMemsetD32Async, dst, value, count, stream);
+
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int32_t), count * sizeof(int32_t), stream, true));
+}
+
 hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
   HIP_INIT_API(hipMemset2D, dst, pitch, value, width, height);
 
-  hip::syncStreams();
-  amd::HostQueue* queue = hip::getNullStream();
-  HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue));
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), pitch * height, nullptr));
 }
 
 hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
                             size_t width, size_t height, hipStream_t stream) {
   HIP_INIT_API(hipMemset2DAsync, dst, pitch, value, width, height, stream);
 
-  amd::HostQueue* queue = hip::getQueue(stream);
-
-  HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue, true));
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), pitch * height, stream, true));
 }
 
-hipError_t ihipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes, hipStream_t stream) {
+hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent) {
+  HIP_INIT_API(hipMemset3D, pitchedDevPtr, value, &extent);
 
-  if (dst == nullptr) {
-    return hipErrorInvalidValue;
-  }
+  void *dst = pitchedDevPtr.ptr;
+  size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
 
-  size_t offset = 0;
-  amd::HostQueue* queue = hip::getQueue(stream);
-  amd::Command::EventWaitList waitList;
-  amd::Memory* memory = getMemoryObject(dst, offset);
-  if (memory != nullptr) {
-    // Device memory
-    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);
-
-    if (command == nullptr) {
-      HIP_RETURN(hipErrorOutOfMemory);
-    }
-
-    command->enqueue();
-    command->awaitCompletion();
-    command->release();
-  } else {
-    // Host alloced memory
-    memset(dst, value, sizeBytes);
-  }
-
-  return hipSuccess;
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), sizeBytes, nullptr));
 }
 
-hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes) {
-  HIP_INIT_API(hipMemsetD8, dst, value, sizeBytes);
+hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent, hipStream_t stream) {
+  HIP_INIT_API(hipMemset3DAsync, pitchedDevPtr, value, &extent, stream);
 
-  HIP_RETURN(ihipMemsetD8(dst, value, sizeBytes, nullptr));
-}
+  void *dst = pitchedDevPtr.ptr;
+  size_t sizeBytes = pitchedDevPtr.pitch * extent.height * extent.depth;
 
-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));
+  HIP_RETURN(ihipMemset(dst, value, sizeof(int8_t), sizeBytes, stream, true));
 }
 
 hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes,