Revert "SWDEV-440866 - [hip-roclr] Adds support to batch memory operations APIs"

This reverts commit bd5d8e9baf.

Reason for revert: hipInfo fails on windows. Updating llvm amd-mainline-closed

Change-Id: I57e1fa1945188b0bc0a799c4f3d540f2b7713003

rocclr/device/blit.hpp

@@ -233,12 +233,6 @@ class BlitManager : public amd::HeapObject {
                              uint64_t mask
   ) const = 0;
 
-  //! Stream batch memory operation
-  virtual bool batchMemOps(const void* paramArray,
-                           size_t paramSize,
-                           uint32_t count
-  ) const = 0;
-
   //! Enables synchronization on blit operations
   void enableSynchronization() { syncOperation_ = true; }
 

rocclr/device/blitcl.cpp

@@ -43,8 +43,6 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
 
   extern void __amd_streamOpsWait(__global uint*, __global ulong*, ulong, ulong, ulong);
 
-  extern void __amd_batchMemOp(__global void*, uint count);
-
   extern void __ockl_dm_init_v1(ulong, ulong, uint, uint);
 
   extern void __ockl_gws_init(uint nwm1, uint rid);
@@ -164,10 +162,6 @@ const char* BlitLinearSourceCode = BLIT_KERNELS(
                                                    ulong4 srcRect, ulong4 dstRect, ulong4 size) {
     __amd_copyBufferRectAligned(src, dst, srcRect, dstRect, size);
   }
-
-  __kernel void __amd_rocclr_batchMemOp(__global void* params, uint count) {
-    __amd_batchMemOp(params, count);
-  }
 );
 
 const char* HipExtraSourceCode = BLIT_KERNELS(
@@ -260,6 +254,7 @@ const char* BlitImageSourceCode = BLIT_KERNELS(
     __amd_copyImageToBuffer(src, dstUInt, dstUShort, dstUChar, srcOrigin, dstOrigin, size, format,
                             pitch);
   }
+
 );
 
 }  // namespace amd::device

rocclr/device/device.hpp

@@ -92,7 +92,6 @@ class SvmMapMemoryCommand;
 class SvmUnmapMemoryCommand;
 class SvmPrefetchAsyncCommand;
 class StreamOperationCommand;
-class BatchMemoryOperationCommand;
 class VirtualMapCommand;
 class ExternalSemaphoreCmd;
 class Isa;
@@ -1309,9 +1308,6 @@ class VirtualDevice : public amd::HeapObject {
     ShouldNotReachHere();
   }
   virtual void submitStreamOperation(amd::StreamOperationCommand& cmd) { ShouldNotReachHere(); }
-  virtual void submitBatchMemoryOperation(amd::BatchMemoryOperationCommand& cmd) {
-    ShouldNotReachHere();
-  }
   virtual void submitVirtualMap(amd::VirtualMapCommand& cmd) { ShouldNotReachHere(); }
 
   virtual address allocKernelArguments(size_t size, size_t alignment) { return nullptr; }

rocclr/device/pal/palblit.hpp

@@ -475,13 +475,6 @@ class KernelBlitManager : public DmaBlitManager {
                         uint number_of_initial_blocks
                         ) const;
 
-  //! Batch memory ops- Submits batch of streamWaits and streamWrite operations.
-  virtual bool batchMemOps(const void* paramArray,
-                           size_t paramSize,
-                           uint32_t count) const {
-    assert(!"Unimplemented");
-    return false;
-  }
  private:
   static constexpr size_t MaxXferBuffers = 2;
   static constexpr uint TransferSplitSize = 3;

rocclr/device/rocm/rocblit.cpp

@@ -2583,38 +2583,6 @@ bool KernelBlitManager::streamOpsWait(device::Memory& memory, uint64_t value, si
   return result;
 }
 
-// ================================================================================================
-bool KernelBlitManager::batchMemOps(const void* paramArray, size_t paramSize,
-                                    uint32_t count) const {
-  amd::ScopedLock k(lockXferOps_);
-  bool result = false;
-  uint blitType = BatchMemOp;
-  size_t dim = 1;
-
-  size_t globalWorkOffset[1] = { 0 };
-  size_t globalWorkSize[1] = { count };
-  size_t localWorkSize[1] = { 1 };
-
-  // Get constant buffer and copy the array of parameters
-  constexpr bool kDirectVa = true;
-  auto constBuf = gpu().allocKernArg((count * paramSize), kCBAlignment);
-  memcpy(constBuf, paramArray, (count * paramSize));
-
-  setArgument(kernels_[blitType], 0, sizeof(cl_mem), constBuf, 0, nullptr, kDirectVa);
-  setArgument(kernels_[blitType], 1, sizeof(cl_mem), &count);
-
-  // Create ND range object for the kernel's execution
-  amd::NDRangeContainer ndrange(dim, globalWorkOffset, globalWorkSize, localWorkSize);
-
-  // Execute the blit
-  address parameters = captureArguments(kernels_[blitType]);
-  result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr);
-  releaseArguments(parameters);
-  synchronize();
-
-  return result;
-}
-
 // ================================================================================================
 bool KernelBlitManager::initHeap(device::Memory* heap_to_initialize, device::Memory* initial_blocks,
                                  uint heap_size, uint number_of_initial_blocks) const {

rocclr/device/rocm/rocblit.hpp

@@ -292,7 +292,6 @@ class KernelBlitManager : public DmaBlitManager {
     BlitCopyBufferRectAligned,
     StreamOpsWrite,
     StreamOpsWait,
-    BatchMemOp,
     Scheduler,
     GwsInit,
     InitHeap,
@@ -520,9 +519,6 @@ class KernelBlitManager : public DmaBlitManager {
                              uint64_t mask
   ) const;
 
-  //! Batch memory ops- Submits batch of streamWaits and streamWrite operations.
-  virtual bool batchMemOps(const void* paramArray, size_t paramSize, uint32_t count) const;
-
   virtual amd::Monitor* lockXfer() const { return &lockXferOps_; }
 
   virtual bool initHeap(device::Memory* heap_to_initialize,
@@ -603,9 +599,9 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = {
   "__amd_rocclr_fillBufferAligned", "__amd_rocclr_fillBufferAligned2D", "__amd_rocclr_copyBuffer",
   "__amd_rocclr_copyBufferAligned", "__amd_rocclr_copyBufferRect",
   "__amd_rocclr_copyBufferRectAligned", "__amd_rocclr_streamOpsWrite", "__amd_rocclr_streamOpsWait",
-  "__amd_rocclr_batchMemOp", "__amd_rocclr_scheduler", "__amd_rocclr_gwsInit",
-  "__amd_rocclr_initHeap", "__amd_rocclr_fillImage", "__amd_rocclr_copyImage",
-  "__amd_rocclr_copyImage1DA", "__amd_rocclr_copyImageToBuffer", "__amd_rocclr_copyBufferToImage"
+  "__amd_rocclr_scheduler", "__amd_rocclr_gwsInit", "__amd_rocclr_initHeap",
+  "__amd_rocclr_fillImage", "__amd_rocclr_copyImage", "__amd_rocclr_copyImage1DA",
+  "__amd_rocclr_copyImageToBuffer", "__amd_rocclr_copyBufferToImage"
 };
 
 inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index,

rocclr/device/rocm/rocvirtual.cpp

@@ -2715,7 +2715,8 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
     else {
     // mask is applied on value before performing
     // the comparision defined by 'condition'
-      bool result = blitMgr().streamOpsWait(*memory, value, offset, sizeBytes, flags, mask);
+      bool result = static_cast<KernelBlitManager&>(blitMgr()).streamOpsWait(*memory, value, offset,
+                                                                              sizeBytes, flags, mask);
       ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Waiting for value: 0x%lx."
               " Flags: 0x%lx mask: 0x%lx", value, flags, mask);
       if (!result) {
@@ -2728,7 +2729,8 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
     // Ensure memory ordering preceding the write
     dispatchBarrierPacket(kBarrierPacketReleaseHeader);
 
-    bool result = blitMgr().streamOpsWrite(*memory, value, offset, sizeBytes);
+    bool result = static_cast<KernelBlitManager&>(blitMgr()).streamOpsWrite(*memory, value,
+                                                                            offset, sizeBytes);
     ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Writing value: 0x%lx", value);
     if (!result) {
       LogError("submitStreamOperation: Write failed!");
@@ -2739,18 +2741,6 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
   profilingEnd(cmd);
 }
 
-// ================================================================================================
-void VirtualGPU::submitBatchMemoryOperation(amd::BatchMemoryOperationCommand& cmd) {
-  // Make sure VirtualGPU has an exclusive access to the resources
-  amd::ScopedLock lock(execution());
-  profilingBegin(cmd);
-
-  bool result = blitMgr().batchMemOps(cmd.getParamPtr(), cmd.paramSize(), cmd.count());
-  if (!result) {
-    LogError("submitBatchMemoryOperation failed!");
-  }
-  profilingEnd(cmd);
-}
 // ================================================================================================
 void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) {
   // Make sure VirtualGPU has an exclusive access to the resources

rocclr/device/rocm/rocvirtual.hpp

@@ -353,7 +353,6 @@ class VirtualGPU : public device::VirtualDevice {
   void flush(amd::Command* list = nullptr, bool wait = false);
   void submitFillMemory(amd::FillMemoryCommand& cmd);
   void submitStreamOperation(amd::StreamOperationCommand& cmd);
-  void submitBatchMemoryOperation(amd::BatchMemoryOperationCommand& cmd);
   void submitVirtualMap(amd::VirtualMapCommand& cmd);
   void submitMigrateMemObjects(amd::MigrateMemObjectsCommand& cmd);
 

rocclr/platform/command.hpp

@@ -845,7 +845,7 @@ class FillMemoryCommand : public OneMemoryArgCommand {
 
 class StreamOperationCommand : public OneMemoryArgCommand {
  private:
-  uint64_t value_;      // !< Value to Wait on or to Write.
+  uint64_t value_;       // !< Value to Wait on or to Write.
   uint64_t mask_;       // !< Mask to be applied on signal value for Wait operation.
   unsigned int flags_;  // !< Flags defining the Wait condition.
   size_t offset_;       // !< Offset into memory for Write
@@ -888,47 +888,6 @@ class StreamOperationCommand : public OneMemoryArgCommand {
   const size_t sizeBytes() const { return sizeBytes_; }
 };
 
-/*! \brief      A batch memory operation command.
- *
- *  \details    Batch operations to synchronize the stream via memory operations
- *              Operations are either 32-bit stream wait or write.
- *              Wait: All the commands issued after stream wait are not executed
- *              until the wait condition is true.
- *              Write: Writes a 32 or 64 bit vaue to the memory using a GPU Blit.
- *              The operations are enqueued in the order they appear in the array.
- */
-
-class BatchMemoryOperationCommand : public Command {
- public:
-  BatchMemoryOperationCommand(HostQueue& queue, cl_command_type cmdType, uint32_t count,
-                              uint32_t flags, EventWaitList& eventWaitList, const void* paramArray,
-                              size_t paramSize)
-    : Command(queue, cmdType, eventWaitList),
-      count_(count),
-      paramArray_(paramArray),
-      flags_(flags),
-      paramSize_(paramSize) {
-    // Sanity check
-    assert(((cmdType == ROCCLR_COMMAND_BATCH_STREAM)) &&
-           "Invalid batch memory operation");
-  }
-
-  virtual void submit(device::VirtualDevice& device) { device.submitBatchMemoryOperation(*this); }
-
-  //! Returns the value
-  const uint64_t count() const { return count_; }
-  //! Return the pointer to the paramList
-  const void* getParamPtr() {return paramArray_; }
-   //! Return the size of a single mem op param in bytes
-  const size_t paramSize() const { return paramSize_; }
-
-  private:
-  uint32_t count_;            // !< The number of operations in the array.
-  uint32_t flags_;            // !< Reserved for future expansion. Must be 0.
-  const void* paramArray_;    // !< Pointer to the array of individual operations
-  size_t paramSize_;          // !< size in bytes of the param array passed
-};
-
 /*! \brief      A generic copy memory command
  *
  *  \details    Used for both buffers and images. Backends are expected
@@ -1899,7 +1858,6 @@ union ComputeCommand {
   CopyMemoryP2PCommand          cmd25;
   SvmPrefetchAsyncCommand       cmd26;
   VirtualMapCommand             cmd27;
-  BatchMemoryOperationCommand   cmd28;
   ComputeCommand() {}
   ~ComputeCommand() {}
 };

rocclr/platform/command_utils.hpp

@@ -24,7 +24,6 @@
 // Dummy command types for Stream Wait and Write commands.
 #define ROCCLR_COMMAND_STREAM_WAIT_VALUE 0x4501
 #define ROCCLR_COMMAND_STREAM_WRITE_VALUE 0x4502
-#define ROCCLR_COMMAND_BATCH_STREAM 0x4503
 
 // Stream Wait Value Conidtions
 #define ROCCLR_STREAM_WAIT_VALUE_GTE 0x0