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

Sourabh Betigeri · agunashe · commit 2ca644cf2255 · 2024-12-02T16:46:12.000-05:00
This reverts commit bd5d8e9. Reason for revert: hipInfo fails on windows. Updating llvm amd-mainline-closed Change-Id: I57e1fa1945188b0bc0a799c4f3d540f2b7713003
diff --git a/hipamd/include/hip/amd_detail/hip_api_trace.hpp b/hipamd/include/hip/amd_detail/hip_api_trace.hpp
@@ -61,7 +61,7 @@
 // - Reset any of the *_STEP_VERSION defines to zero if the corresponding *_MAJOR_VERSION increases
 #define HIP_API_TABLE_STEP_VERSION 0
 #define HIP_COMPILER_API_TABLE_STEP_VERSION 0
-#define HIP_RUNTIME_API_TABLE_STEP_VERSION 7
+#define HIP_RUNTIME_API_TABLE_STEP_VERSION 6
 
 // HIP API interface
 typedef hipError_t (*t___hipPopCallConfiguration)(dim3* gridDim, dim3* blockDim, size_t* sharedMem,
@@ -722,8 +722,6 @@ typedef hipError_t (*t_hipStreamWriteValue32)(hipStream_t stream, void* ptr, uin
                                               unsigned int flags);
 typedef hipError_t (*t_hipStreamWriteValue64)(hipStream_t stream, void* ptr, uint64_t value,
                                               unsigned int flags);
-typedef hipError_t (*t_hipStreamBatchMemOp)(hipStream_t stream, unsigned int count,
-    hipStreamBatchMemOpParams* paramArray, unsigned int flags);
 typedef hipError_t (*t_hipTexObjectCreate)(hipTextureObject_t* pTexObject,
                                            const HIP_RESOURCE_DESC* pResDesc,
                                            const HIP_TEXTURE_DESC* pTexDesc,
@@ -1521,9 +1519,6 @@ struct HipDispatchTable {
   // HIP_RUNTIME_API_TABLE_STEP_VERSION == 6
   t_hipDeviceGetTexture1DLinearMaxWidth hipDeviceGetTexture1DLinearMaxWidth_fn;
 
-  // HIP_RUNTIME_API_TABLE_STEP_VERSION == 7
-  t_hipStreamBatchMemOp hipStreamBatchMemOp_fn;
-
   // DO NOT EDIT ABOVE!
   // HIP_RUNTIME_API_TABLE_STEP_VERSION == 7
 
diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h
@@ -425,8 +425,7 @@ enum hip_api_id_t {
   HIP_API_ID_hipMemcpyHtoAAsync = 405,
   HIP_API_ID_hipSetValidDevices = 406,
   HIP_API_ID_hipExtHostAlloc = 407,
-  HIP_API_ID_hipStreamBatchMemOp = 408,
-  HIP_API_ID_LAST = 408,
+  HIP_API_ID_LAST = 407,
 
   HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
   HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties),
@@ -546,6 +545,7 @@ static inline const char* hip_api_name(const uint32_t id) {
     case HIP_API_ID_hipEventQuery: return "hipEventQuery";
     case HIP_API_ID_hipEventRecord: return "hipEventRecord";
     case HIP_API_ID_hipEventSynchronize: return "hipEventSynchronize";
+    case HIP_API_ID_hipExtGetLastError: return "hipExtGetLastError";
     case HIP_API_ID_hipExtGetLinkTypeAndHopCount: return "hipExtGetLinkTypeAndHopCount";
     case HIP_API_ID_hipExtLaunchKernel: return "hipExtLaunchKernel";
     case HIP_API_ID_hipExtLaunchMultiKernelMultiDevice: return "hipExtLaunchMultiKernelMultiDevice";
@@ -859,8 +859,6 @@ static inline const char* hip_api_name(const uint32_t id) {
     case HIP_API_ID_hipUserObjectRelease: return "hipUserObjectRelease";
     case HIP_API_ID_hipUserObjectRetain: return "hipUserObjectRetain";
     case HIP_API_ID_hipWaitExternalSemaphoresAsync: return "hipWaitExternalSemaphoresAsync";
-    case HIP_API_ID_hipExtGetLastError: return "hipExtGetLastError";
-    case HIP_API_ID_hipStreamBatchMemOp: return "hipStreamBatchMemOp";
   };
   return "unknown";
 };
@@ -1264,7 +1262,6 @@ static inline uint32_t hipApiIdByName(const char* name) {
   if (strcmp("hipUserObjectRelease", name) == 0) return HIP_API_ID_hipUserObjectRelease;
   if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain;
   if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync;
-  if (strcmp("hipStreamBatchMemOp", name) == 0) return HIP_API_ID_hipStreamBatchMemOp;
   return HIP_API_ID_NONE;
 }
 
@@ -3626,13 +3623,6 @@ typedef struct hip_api_data_s {
       unsigned int numExtSems;
       hipStream_t stream;
     } hipWaitExternalSemaphoresAsync;
-    struct {
-      hipStream_t stream;
-      unsigned int count;
-      hipStreamBatchMemOpParams* paramArray;
-      hipStreamBatchMemOpParams paramArray__val;
-      unsigned int flags;
-    } hipStreamBatchMemOp;
   } args;
   uint64_t *phase_data;
 } hip_api_data_t;
@@ -5902,15 +5892,6 @@ typedef struct hip_api_data_s {
   cb_data.args.hipStreamWriteValue64.value = (uint64_t)value; \
   cb_data.args.hipStreamWriteValue64.flags = (unsigned int)flags; \
 };
-
-// hipStreamBatchMemOp[('hipStream_t', 'stream'), ('unsigned int', 'count'),
-// ('hipStreamBatchMemOpParams*', 'paramArray'), ('unsigned int', 'flags')]
-#define INIT_hipStreamBatchMemOp_CB_ARGS_DATA(cb_data) { \
-  cb_data.args.hipStreamBatchMemOp.stream = (hipStream_t)stream; \
-  cb_data.args.hipStreamBatchMemOp.count = (unsigned int)count; \
-  cb_data.args.hipStreamBatchMemOp.paramArray= (hipStreamBatchMemOpParams*)paramArray; \
-  cb_data.args.hipStreamBatchMemOp.flags = (unsigned int)flags; \
-};
 // hipTexRefGetAddress[('hipDeviceptr_t*', 'dev_ptr'), ('const textureReference*', 'texRef')]
 #define INIT_hipTexRefGetAddress_CB_ARGS_DATA(cb_data) { \
   cb_data.args.hipTexRefGetAddress.dev_ptr = (hipDeviceptr_t*)dptr; \
@@ -7546,11 +7527,6 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
 // hipStreamWriteValue64[('hipStream_t', 'stream'), ('void*', 'ptr'), ('uint64_t', 'value'), ('unsigned int', 'flags')]
     case HIP_API_ID_hipStreamWriteValue64:
       break;
-// hipStreamBatchMemOp[('hipStream_t', 'stream'), ('unsigned int', 'count'),
-// ('hipStreamBatchMemOpParams*', 'paramArray'), ('unsigned int', 'flags')]
-    case HIP_API_ID_hipStreamBatchMemOp:
-      if (data->args.hipStreamBatchMemOp.paramArray) data->args.hipStreamBatchMemOp.paramArray__val = *(data->args.hipStreamBatchMemOp.paramArray);
-      break;
 // hipTexRefGetAddress[('hipDeviceptr_t*', 'dev_ptr'), ('const textureReference*', 'texRef')]
     case HIP_API_ID_hipTexRefGetAddress:
       if (data->args.hipTexRefGetAddress.dev_ptr) data->args.hipTexRefGetAddress.dev_ptr__val = *(data->args.hipTexRefGetAddress.dev_ptr);
@@ -10638,15 +10614,6 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
       oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamWriteValue64.flags);
       oss << ")";
     break;
-    case HIP_API_ID_hipStreamBatchMemOp:
-      oss << "hipStreamBatchMemOp(";
-      oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.stream);
-      oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.count);
-      if (data->args.hipStreamBatchMemOp.paramArray == NULL) oss << ", paramArray=NULL";
-      else { oss << ", paramArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.paramArray__val); }
-      oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamBatchMemOp.flags);
-      oss << ")";
-    break;
     case HIP_API_ID_hipTexRefGetAddress:
       oss << "hipTexRefGetAddress(";
       if (data->args.hipTexRefGetAddress.dev_ptr == NULL) oss << "dev_ptr=NULL";
diff --git a/hipamd/src/CMakeLists.txt b/hipamd/src/CMakeLists.txt
@@ -299,7 +299,7 @@ if(WIN32 OR NOT BUILD_SHARED_LIBS)
     # rocprofiler-register is not support on Windows
     set(HIP_ENABLE_ROCPROFILER_REGISTER OFF)
 else()
-  option(HIP_ENABLE_ROCPROFILER_REGISTER "Enable rocprofiler-register support" ON)
+    option(HIP_ENABLE_ROCPROFILER_REGISTER "Enable rocprofiler-register support" ON)
 endif()
 
 if(HIP_ENABLE_ROCPROFILER_REGISTER)
diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def
@@ -480,4 +480,3 @@ hipGraphExecNodeSetParams
 hipDrvGraphMemcpyNodeSetParams
 hipDrvGraphMemcpyNodeGetParams
 hipExtHostAlloc
-hipStreamBatchMemOp
diff --git a/hipamd/src/hip_api_trace.cpp b/hipamd/src/hip_api_trace.cpp
@@ -602,8 +602,6 @@ hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, uint64_t value, u
                                 uint64_t mask);
 hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags);
 hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags);
-hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count, hipStreamBatchMemOpParams* paramArray,
-                               unsigned int flags);
 hipError_t hipTexObjectCreate(hipTextureObject_t* pTexObject, const HIP_RESOURCE_DESC* pResDesc,
                               const HIP_TEXTURE_DESC* pTexDesc,
                               const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
@@ -1199,7 +1197,6 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
   ptrDispatchTable->hipStreamWaitValue64_fn = hip::hipStreamWaitValue64;
   ptrDispatchTable->hipStreamWriteValue32_fn = hip::hipStreamWriteValue32;
   ptrDispatchTable->hipStreamWriteValue64_fn = hip::hipStreamWriteValue64;
-  ptrDispatchTable->hipStreamBatchMemOp_fn = hip::hipStreamBatchMemOp;
   ptrDispatchTable->hipTexObjectCreate_fn = hip::hipTexObjectCreate;
   ptrDispatchTable->hipTexObjectDestroy_fn = hip::hipTexObjectDestroy;
   ptrDispatchTable->hipTexObjectGetResourceDesc_fn = hip::hipTexObjectGetResourceDesc;
@@ -1890,19 +1887,16 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipDrvGraphMemcpyNodeSetParams_fn, 460)
 HIP_ENFORCE_ABI(HipDispatchTable, hipExtHostAlloc_fn, 461)
 // HIP_RUNTIME_API_TABLE_STEP_VERSION == 6
 HIP_ENFORCE_ABI(HipDispatchTable, hipDeviceGetTexture1DLinearMaxWidth_fn, 462)
-// HIP_RUNTIME_API_TABLE_STEP_VERSION == 7
-HIP_ENFORCE_ABI(HipDispatchTable, hipStreamBatchMemOp_fn, 463);
-
 
 // if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below
 // will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.:
 //
 //  HIP_ENFORCE_ABI(<table>, <functor>, 8)
 //
 //  HIP_ENFORCE_ABI_VERSIONING(<table>, 9) <- 8 + 1 = 9
-HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 464)
+HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 463)
 
-static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 7,
+static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 6,
               "If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
               "pointers and then update this check so it is true");
 #endif
diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in
@@ -583,7 +583,6 @@ local:
 hip_6.3 {
 global:
     hipExtHostAlloc;
-    hipStreamBatchMemOp;
 local:
     *;
 } hip_6.2;
diff --git a/hipamd/src/hip_stream_ops.cpp b/hipamd/src/hip_stream_ops.cpp
@@ -23,36 +23,8 @@
 #include "platform/command_utils.hpp"
 
 namespace hip {
-hipError_t ihipBatchMemOperation(hipStream_t stream, cl_command_type cmdType, unsigned int count,
-                                 hipStreamBatchMemOpParams* paramArray, unsigned int flags) {
-  if (paramArray == nullptr || flags != 0 || count > 256) {
-    return hipErrorInvalidValue;
-  }
-
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
-
-  hip::Stream* hip_stream = hip::getStream(stream);
-  amd::Command::EventWaitList waitList;
-
-  amd::BatchMemoryOperationCommand* command =
-    new amd::BatchMemoryOperationCommand(*hip_stream, cmdType, count,
-                                         flags, waitList, paramArray,
-                                         sizeof(hipStreamBatchMemOpParams));
-
-  if (command == nullptr) {
-    return hipErrorOutOfMemory;
-  }
-  command->enqueue();
-  command->release();
-  HIP_RETURN(hipSuccess);
-}
-
-
 hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr,
-                               uint64_t value, uint64_t mask, unsigned int flags,
-                               size_t sizeBytes) {
+                               uint64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes) {
   size_t offset = 0;
   unsigned int outFlags = 0;
 
@@ -164,15 +136,4 @@ hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value,
       0,  // flags un-used for now set it to 0
       sizeof(uint64_t)));
 }
-
-hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count,
-                               hipStreamBatchMemOpParams* paramArray, unsigned int flags) {
-  HIP_INIT_API(hipStreamBatchMemOp, count, paramArray, flags);
-  HIP_RETURN_DURATION(ihipBatchMemOperation(
-                      stream,
-                      ROCCLR_COMMAND_BATCH_STREAM,
-                      count,
-                      paramArray,
-                      flags));
-  }
 }  // namespace hip
diff --git a/hipamd/src/hip_table_interface.cpp b/hipamd/src/hip_table_interface.cpp
@@ -1413,10 +1413,6 @@ hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value,
                                  unsigned int flags) {
   return hip::GetHipDispatchTable()->hipStreamWriteValue64_fn(stream, ptr, value, flags);
 }
-hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count,
-                               hipStreamBatchMemOpParams* paramArray, unsigned int flags) {
-  return hip::GetHipDispatchTable()->hipStreamBatchMemOp_fn(stream, count, paramArray, flags);
-}
 hipError_t hipTexObjectCreate(hipTextureObject_t* pTexObject, const HIP_RESOURCE_DESC* pResDesc,
                               const HIP_TEXTURE_DESC* pTexDesc,
                               const HIP_RESOURCE_VIEW_DESC* pResViewDesc) {
diff --git a/rocclr/device/blit.hpp b/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; }
 
diff --git a/rocclr/device/blitcl.cpp b/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
diff --git a/rocclr/device/device.hpp b/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; }
diff --git a/rocclr/device/pal/palblit.hpp b/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;
diff --git a/rocclr/device/rocm/rocblit.cpp b/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 {
diff --git a/rocclr/device/rocm/rocblit.hpp b/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,
diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp
diff --git a/rocclr/platform/command.hpp b/rocclr/platform/command.hpp
diff --git a/rocclr/platform/command_utils.hpp b/rocclr/platform/command_utils.hpp
