//
// Copyright (c) 2009 Advanced Micro Devices, Inc. All rights reserved.
//
#include "cl_common.hpp"
#include "platform/command.hpp"
#include "platform/kernel.hpp"
#include "platform/program.hpp"

/*! \brief Helper function to validate SVM allocation flags
 *
 *  \return true if flags are valid, otherwise - false
 */
static bool validateSvmFlags(cl_svm_mem_flags flags) {
  if (!flags) {
    // coarse-grained allocation
    return true;
  }
  const cl_svm_mem_flags rwFlags = CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY;
  const cl_svm_mem_flags setFlags =
      flags & (rwFlags | CL_MEM_SVM_ATOMICS | CL_MEM_SVM_FINE_GRAIN_BUFFER);
  if (flags != setFlags) {
    // invalid flags value
    return false;
  }

  if (amd::countBitsSet(flags & rwFlags) > 1) {
    // contradictory R/W flags
    return false;
  }

  if ((flags & CL_MEM_SVM_ATOMICS) && !(flags & CL_MEM_SVM_FINE_GRAIN_BUFFER)) {
    return false;
  }

  return true;
}

/*! \brief Helper function to validate cl_map_flags
 *
 *  \return true if flags are valid, otherwise - false
 */
static bool validateMapFlags(cl_map_flags flags) {
  const cl_map_flags maxFlag = CL_MAP_WRITE_INVALIDATE_REGION;
  if (flags >= (maxFlag << 1)) {
    // at least one flag is out-of-range
    return false;
  } else if ((flags & CL_MAP_WRITE_INVALIDATE_REGION) && (flags & (CL_MAP_READ | CL_MAP_WRITE))) {
    // CL_MAP_READ or CL_MAP_WRITE and CL_MAP_WRITE_INVALIDATE_REGION are
    // mutually exclusive.
    return false;
  }
  return true;
}

/*! \addtogroup API
 *  @{
 *
 *  \addtogroup SVM
 *  @{
 *
 */

/*! \brief Allocate a shared virtual memory buffer that can be shared by the
 * host and all devices in an OpenCL context.
 *
 *  \param context is a valid OpenCL context used to create the SVM buffer.
 *
 *  \param flags is a bit-field that is used to specify allocation and usage
 *  information. If CL_MEM_SVM_FINE_GRAIN_BUFFER is not specified, the
 *  buffer is created as a coarse grained SVM allocation. Similarly, if
 *  CL_MEM_SVM_ATOMICS is not specified, the buffer is created without
 *  support for SVM atomic operations.
 *
 *  \param size is the size in bytes of the SVM buffer to be allocated.
 *
 *  \param alignment is the minimum alignment in bytes that is required for the
 *  newly created buffer?s memory region. It must be a power of two up to the
 *  largest data type supported by the OpenCL device. For the full profile, the
 *  largest data type is long16. For the embedded profile, it is long16 if the
 *  device supports 64-bit integers; otherwise it is int16. If alignment is 0, a
 *  default alignment will be used that is equal to the size of largest data
 *  type supported by the OpenCL implementation.
 *
 *  \return A valid non-NULL shared virtual memory address if the SVM buffer
 *  is successfully allocated. Otherwise, like malloc, it returns a NULL pointer
 *  value. clSVMAlloc will fail if
 *  - \a context is not a valid context.
 *  - \a flags does not contain CL_MEM_SVM_FINE_GRAIN_BUFFER but does
 *  contain CL_MEM_SVM_ATOMICS.
 *  - Values specified in \a flags do not follow rules for that particular type.
 *  - CL_MEM_SVM_FINE_GRAIN_BUFFER or CL_MEM_SVM_ATOMICS is specified
 *  in \a flags and these are not supported by at least one device in
 *  \a context.
 *  - The values specified in \a flags are not valid.
 *  - \a size is 0 or > CL_DEVICE_MAX_MEM_ALLOC_SIZE value for any device in
 *  \a context.
 *  - \a alignment is not a power of two or the OpenCL implementation cannot
 *  support the specified alignment for at least one device in \a context.
 *  - There was a failure to allocate resources.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY_RET_NOERRCODE(void*, clSVMAlloc, (cl_context context, cl_svm_mem_flags flags,
                                                size_t size, unsigned int alignment)) {
  if (!is_valid(context)) {
    LogWarning("invalid parameter \"context\"");
    return NULL;
  }

  if (size == 0) {
    LogWarning("invalid parameter \"size = 0\"");
    return NULL;
  }

  if (!validateSvmFlags(flags)) {
    LogWarning("invalid parameter \"flags\"");
    return NULL;
  }

  if (!amd::isPowerOfTwo(alignment)) {
    LogWarning("invalid parameter \"alignment\"");
    return NULL;
  }

  const std::vector<amd::Device*>& devices = as_amd(context)->svmDevices();
  bool sizePass = false;
  cl_device_svm_capabilities combinedSvmCapabilities = 0;
  const cl_uint hostAddressBits = LP64_SWITCH(32, 64);
  cl_uint minContextAlignment = std::numeric_limits<uint>::max();
  for (const auto& it : devices) {
    cl_device_svm_capabilities svmCapabilities = it->info().svmCapabilities_;
    if (svmCapabilities == 0) {
      continue;
    }
    combinedSvmCapabilities |= svmCapabilities;

    if (it->info().maxMemAllocSize_ >= size) {
      sizePass = true;
    }

    if (it->info().addressBits_ < hostAddressBits) {
      LogWarning("address mode mismatch between host and device");
      return NULL;
    }

    // maximum alignment for a device is given in bits.
    cl_uint baseAlignment = it->info().memBaseAddrAlign_ >> 3;
    if (alignment > baseAlignment) {
      LogWarning("invalid parameter \"alignment\"");
      return NULL;
    }

    minContextAlignment = std::min(minContextAlignment, baseAlignment);
  }
  if ((flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) &&
      !(combinedSvmCapabilities & CL_DEVICE_SVM_FINE_GRAIN_BUFFER)) {
    LogWarning("No device in context supports SVM fine grained buffers");
    return NULL;
  }
  if ((flags & CL_MEM_SVM_ATOMICS) && !(combinedSvmCapabilities & CL_DEVICE_SVM_ATOMICS)) {
    LogWarning("No device in context supports SVM atomics");
    return NULL;
  }
  if (!sizePass) {
    LogWarning("invalid parameter \"size\"");
    return NULL;
  }

  // if alignment not specified, use largest data type alignment supported
  if (alignment == 0) {
    alignment = minContextAlignment;
    ClPrint(amd::LOG_INFO, amd::LOG_API, "Assumed alignment %d\n", alignment);
  }

  amd::Context& amdContext = *as_amd(context);
  return amd::SvmBuffer::malloc(amdContext, flags, size, alignment);
}
RUNTIME_EXIT

/*! \brief Free a shared virtual memory buffer allocated using clSVMAlloc.
 *
 *  \param context is a valid OpenCL context used to create the SVM buffer.
 *
 *  \param svm_pointer must be the value returned by a call to clSVMAlloc. If a
 *  NULL pointer is passed in \a svm_pointer, no action occurs.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY_VOID(void, clSVMFree, (cl_context context, void* svm_pointer)) {
  if (!is_valid(context)) {
    LogWarning("invalid parameter \"context\"");
    return;
  }

  if (svm_pointer == NULL) {
    return;
  }

  amd::Context& amdContext = *as_amd(context);
  amd::SvmBuffer::free(amdContext, svm_pointer);
}
RUNTIME_EXIT

/*! \brief enqueues a command to free shared virtual memory allocated using
 *  clSVMAlloc or a shared system memory pointer.
 *
 *  \param command_queue is a valid host command-queue.
 *
 *  \param num_svm_pointers specifies the number of elements in \a svm_pointers.
 *
 *  \param svm_pointers is a list of shared virtual memory pointers to
 *  be freed. Each pointer in \a svm_pointers that was allocated using SVMAlloc
 *  must have been allocated from the same context from which \a command_queue
 *  was created. The memory associated with \a svm_pointers can be reused or
 *  freed after the function returns.
 *
 *  \param pfn_free_func specifies the callback function to be called to free
 *  the SVM pointers. \a pfn_free_func takes four arguments: \a queue which is
 *  the command queue in which clEnqueueSVMFree was enqueued, the count and list
 *  of SVM pointers to free and \a user_data which is a pointer to user
 *  specified data. If \a pfn_free_func is NULL, all the pointers specified in
 *  \a svm_pointers array must be allocated using clSVMAlloc. \a pfn_free_func
 *  must however be a valid callback function if any SVM pointer to be freed is
 *  a shared system memory pointer i.e. not allocated using clSVMAlloc.
 *
 *  \param user_data will be passed as the user_data argument when
 *  \a pfn_free_func is called. \a user_data can be NULL.
 *
 *  \param even_wait_list specifies the  events that need to complete before
 *  this particular command can be executed. If \a event_wait_list is NULL, then
 *  this particular command does not wait on any event to complete. If
 *  \a event_wait_list is NULL, \a num_events_in_wait_list must be 0. If
 *  \a event_wait_list is not NULL, the list of events pointed to by
 *  \a event_wait_list must be valid and \a num_events_in_wait_list must be
 *  greater than 0. The events specified in \a event_wait_list act as
 *  synchronization points. The context associated with events in
 *  \a event_wait_list and \a command_queue must be the same. The memory
 *  associated with \a event_wait_list can be reused or freed after the function
 *  returns.
 *
 *  \param num_events_in_wait_list specifies the number of elements in
 *  \a even_wait_list
 *
 *  \param event returns an event object that identifies this particular command
 *  and can be used to query or queue a wait for this particular command to
 *  complete. \a event can be NULL in which case it will not be possible for the
 *  application to query the status of this command or queue a wait for this
 *  command to complete. If the \a event_wait_list and the \a event arguments
 *  are not NULL, the \a event argument should not refer to an element of the
 *  \a event_wait_list array.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid host
 *    command-queue
 *  - CL_INVALID_VALUE if \a num_svm_pointers is 0 or if \a svm_pointers is
 *    NULL or if any of the pointers specified in \a svm_pointers array is NULL
 *  - CL_INVALID_CONTEXT if context associated with \a command_queue and
 *    events in \a event_wait_list are not the same
 *  - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
 *    \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
 *    \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
 *    are not valid events.
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clEnqueueSVMFree,
              (cl_command_queue command_queue, cl_uint num_svm_pointers, void* svm_pointers[],
               void(CL_CALLBACK* pfn_free_func)(cl_command_queue queue, cl_uint num_svm_pointers,
                                                void* svm_pointers[], void* user_data),
               void* user_data, cl_uint num_events_in_wait_list, const cl_event* event_wait_list,
               cl_event* event)) {
  if (!is_valid(command_queue)) {
    return CL_INVALID_COMMAND_QUEUE;
  }

  if (num_svm_pointers == 0) {
    LogWarning("invalid parameter \"num_svm_pointers = 0\"");
    return CL_INVALID_VALUE;
  }

  if (svm_pointers == NULL) {
    LogWarning("invalid parameter \"svm_pointers = NULL\"");
    return CL_INVALID_VALUE;
  }

  //!@todo why are NULL pointers disallowed here but not in clSVMFree?
  for (cl_uint i = 0; i < num_svm_pointers; i++) {
    if (svm_pointers[i] == NULL) {
      LogWarning("Null pointers are not allowed");
      return CL_INVALID_VALUE;
    }
  }

  //!@todo what if the callback is NULL but \a user_data is not?

  amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
  if (NULL == queue) {
    return CL_INVALID_COMMAND_QUEUE;
  }
  amd::HostQueue& hostQueue = *queue;

  amd::Command::EventWaitList eventWaitList;
  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
                                       event_wait_list);
  if (err != CL_SUCCESS) {
    return err;
  }

  amd::Command* command = new amd::SvmFreeMemoryCommand(hostQueue, eventWaitList, num_svm_pointers,
                                                        svm_pointers, pfn_free_func, user_data);

  if (command == NULL) {
    return CL_OUT_OF_HOST_MEMORY;
  }

  command->enqueue();

  *not_null(event) = as_cl(&command->event());
  if (event == NULL) {
    command->release();
  }

  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief enqueues a command to do a memcpy operation.
 *
 *  \param command_queue refers to the host command-queue in which the read/
 *  write commands will be queued.
 *
 *  \param blocking_copy indicates if the copy operation is blocking or
 *  non-blocking. If \a blocking_copy is CL_TRUE i.e. the copy command is
 *  blocking, clEnqueueSVMMemcpy does not return until the buffer data has been
 *  copied into memory pointed to by \a dst_ptr. If \a blocking_copy is CL_FALSE
 *  i.e. the copy command is non-blocking, clEnqueueSVMMemcpy queues a
 *  non-blocking copy command and returns. The contents of the buffer that
 *  \a dst_ptr point to cannot be used until the copy command has completed.
 *  The \a event argument returns an event object which can be used to query the
 *  execution status of the read command. When the copy command has completed,
 *  the contents of the buffer that \a dst_ptr points to can be used by the
 *  application.
 *
 *  \param dst_ptr is the pointer to a memory region where data is copied to.
 *
 *  \param src_ptr is the pointer to a memory region where data is copied from.
 *  If \a dst_ptr and/or \a src_ptr are allocated using clSVMAlloc then they
 *  must be allocated from the same context from which \a command_queue was
 *  created. Otherwise the behavior is undefined.
 *
 *  \param size is the size in bytes of data being copied.
 *
 *  \param even_wait_list specifies the  events that need to complete before
 *  this particular command can be executed. If \a event_wait_list is NULL, then
 *  this particular command does not wait on any event to complete. If
 *  \a event_wait_list is NULL, \a num_events_in_wait_list must be 0. If
 *  \a event_wait_list is not NULL, the list of events pointed to by
 *  \a event_wait_list must be valid and \a num_events_in_wait_list must be
 *  greater than 0. The events specified in \a event_wait_list act as
 *  synchronization points. The context associated with events in
 *  \a event_wait_list and \a command_queue must be the same. The memory
 *  associated with \a event_wait_list can be reused or freed after the function
 *  returns.
 *
 *  \param num_events_in_wait_list specifies the number of elements in
 *  \a even_wait_list
 *
 *  \param event returns an event object that identifies this particular command
 *  and can be used to query or queue a wait for this particular command to
 *  complete. \a event can be NULL in which case it will not be possible for the
 *  application to query the status of this command or queue a wait for this
 *  command to complete. If the \a event_wait_list and the \a event arguments
 *  are not NULL, the \a event argument should not refer to an element of the
 *  \a event_wait_list array.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid host
 *    command-queue
 *  - CL_INVALID_CONTEXT if the context associated with \a command_queue and
 *    events in \a event_wait_list are not the same
 *  - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
 *    \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
 *    \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
 *    are not valid events.
 *  - CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST if the operation is
 *    blocking and the execution status of any of the events in
 *    \a event_wait_list is a negative integer value.
 *  - CL_INVALID_VALUE if \a dst_ptr or \a src_ptr are NULL.
 *  - CL_INVALID_VALUE if \a size is 0.
 *  - CL_MEM_COPY_OVERLAP if the values specified for \a dst_ptr, \a src_ptr
 *    and \a size result in an overlapping copy.
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clEnqueueSVMMemcpy,
              (cl_command_queue command_queue, cl_bool blocking_copy, void* dst_ptr,
               const void* src_ptr, size_t size, cl_uint num_events_in_wait_list,
               const cl_event* event_wait_list, cl_event* event)) {
  if (!is_valid(command_queue)) {
    return CL_INVALID_COMMAND_QUEUE;
  }

  if (dst_ptr == NULL || src_ptr == NULL) {
    return CL_INVALID_VALUE;
  }

  if (size == 0) {
    return CL_INVALID_VALUE;
  }

  char* dst = reinterpret_cast<char*>(dst_ptr);
  const char* src = reinterpret_cast<const char*>(src_ptr);
  if ((dst > src - size) && (dst < src + size)) {
    return CL_MEM_COPY_OVERLAP;
  }

  amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
  if (NULL == queue) {
    return CL_INVALID_COMMAND_QUEUE;
  }
  amd::HostQueue& hostQueue = *queue;

  amd::Command::EventWaitList eventWaitList;
  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
                                       event_wait_list);
  if (err != CL_SUCCESS) {
    return err;
  }

  amd::Command* command =
      new amd::SvmCopyMemoryCommand(hostQueue, eventWaitList, dst_ptr, src_ptr, size);

  if (command == NULL) {
    return CL_OUT_OF_HOST_MEMORY;
  }

  command->enqueue();

  if (blocking_copy) {
    command->awaitCompletion();
  }

  *not_null(event) = as_cl(&command->event());
  if (event == NULL) {
    command->release();
  }

  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief enqueues a command to fill a region in memory with a pattern of a
 * given pattern size.
 *
 *  \param command_queue refers to the host command-queue in which the fill
 *  command will be queued. The OpenCL context associated with \a command_queue
 *  and SVM pointer referred to by \a svm_ptr must be the same..
 *
 *  \param svm_ptr is a pointer to a memory region that will be filled with
 *  \a pattern. It must be aligned to \a pattern_size bytes. If \a svm_ptr is
 *  allocated using clSVMAlloc then it must be allocated from the same context
 *  from which \a command_queue was created. Otherwise the behavior is
 *  undefined.
 *
 *  \a pattern is a pointer to the data pattern of size \a pattern_size in
 *  bytes. \a pattern will be used to fill a region in buffer starting at
 *  \a svm_ptr and is \a size bytes in size. The data pattern must be a scalar
 *  or vector integer or floating-point data type supported by OpenCL. For
 *  example, if region pointed to by \a svm_ptr is to be filled with a pattern
 *  of float4 values, then \a pattern will be a pointer to a cl_float4 value
 *  and \a pattern_size will be sizeof(cl_float4). The maximum value of
 *  \a pattern_size is the size of the largest integer or floating-point vector
 *  data type supported by the OpenCL device. The memory associated with
 *  \a pattern can be reused or freed after the function returns.
 *
 *  \param size is the size in bytes of region being filled starting with
 *  \a svm_ptr and must be a multiple of \a pattern_size.
 *
 *  \param even_wait_list specifies the  events that need to complete before
 *  this particular command can be executed. If \a event_wait_list is NULL, then
 *  this particular command does not wait on any event to complete. If
 *  \a event_wait_list is NULL, \a num_events_in_wait_list must be 0. If
 *  \a event_wait_list is not NULL, the list of events pointed to by
 *  \a event_wait_list must be valid and \a num_events_in_wait_list must be
 *  greater than 0. The events specified in \a event_wait_list act as
 *  synchronization points. The context associated with events in
 *  \a event_wait_list and \a command_queue must be the same. The memory
 *  associated with \a event_wait_list can be reused or freed after the function
 *  returns.
 *
 *  \param num_events_in_wait_list specifies the number of elements in
 *  \a even_wait_list
 *
 *  \param event returns an event object that identifies this particular command
 *  and can be used to query or queue a wait for this particular command to
 *  complete. \a event can be NULL in which case it will not be possible for the
 *  application to query the status of this command or queue a wait for this
 *  command to complete. clEnqueueBarrierWithWaitList can be used instead. If
 *  the \a event_wait_list and the \a event arguments are not NULL, the \a event
 *  argument should not refer to an element of the \a event_wait_list array.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid host
 *    command-queue
 *  - CL_INVALID_CONTEXT if context associated with \a command_queue and
 *    events in \a event_wait_list are not the same
 *  - CL_INVALID_VALUE if \a svm_ptr is NULL.
 *  - CL_INVALID_VALUE if \a svm_ptr is not aligned to \a pattern_size bytes.
 *  - CL_INVALID_VALUE if \a pattern is NULL or if \a pattern_size is 0 or if
 *    \a pattern_size is not one of {1, 2, 4, 8, 16, 32, 64, 128}.
 *  - CL_INVALID_VALUE if \a size is 0 or is not a multiple of \a pattern_size.
 *  - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
 *    \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
 *    \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
 *    are not valid events.
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clEnqueueSVMMemFill,
              (cl_command_queue command_queue, void* svm_ptr, const void* pattern,
               size_t pattern_size, size_t size, cl_uint num_events_in_wait_list,
               const cl_event* event_wait_list, cl_event* event)) {
  if (!is_valid(command_queue)) {
    return CL_INVALID_COMMAND_QUEUE;
  }

  if (svm_ptr == NULL) {
    return CL_INVALID_VALUE;
  }

  char* dst = reinterpret_cast<char*>(svm_ptr);
  if (!amd::isMultipleOf(dst, pattern_size)) {
    return CL_INVALID_VALUE;
  }

  if (pattern == NULL) {
    return CL_INVALID_VALUE;
  }

  if (!amd::isPowerOfTwo(pattern_size) || pattern_size == 0 ||
      pattern_size > amd::FillMemoryCommand::MaxFillPatterSize) {
    return CL_INVALID_VALUE;
  }

  if (size == 0 || !amd::isMultipleOf(size, pattern_size)) {
    return CL_INVALID_VALUE;
  }

  amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
  if (NULL == queue) {
    return CL_INVALID_COMMAND_QUEUE;
  }
  amd::HostQueue& hostQueue = *queue;

  amd::Command::EventWaitList eventWaitList;
  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
                                       event_wait_list);
  if (err != CL_SUCCESS) {
    return err;
  }

  amd::Command* command =
      new amd::SvmFillMemoryCommand(hostQueue, eventWaitList, svm_ptr, pattern, pattern_size, size);

  if (command == NULL) {
    return CL_OUT_OF_HOST_MEMORY;
  }

  command->enqueue();

  *not_null(event) = as_cl(&command->event());
  if (event == NULL) {
    command->release();
  }

  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief enqueues a command that will allow the host to update a region of a
 *  SVM buffer
 *
 *  \param command_queue is a valid host command-queue.
 *
 *  \param blocking_map indicates if the map operation is blocking or
 *  non-blocking. If \a blocking_map is CL_TRUE, clEnqueueSVMMap does not return
 *  until the application can access the contents of the SVM region specified by
 *  \a svm_ptr and \a size on the host. If blocking_map is CL_FALSE i.e. map
 *  operation is non-blocking, the region specified by \a svm_ptr and \a size
 *  cannot be used until the map command has completed. The \a event argument
 *  returns an event object which can be used to query the execution status of
 *  the map command. When the map command is completed, the application can
 *  access the contents of the region specified by \a svm_ptr and \a size.
 *
 *  \param maps_flag is a valid cl_map_flags flag.
 *
 *  \param svm_ptr is a pointer to a memory region that will be updated by the
 *   host. If \a svm_ptr is allocated using clSVMAlloc then it must be allocated
 *   from the same context from which \a command_queue was created. Otherwise
 *   the behavior is undefined.
 *
 *  \param size is the size in bytes of the memory region that will be updated
 *  by the host.
 *
 *  \param even_wait_list specifies the  events that need to complete before
 *  this particular command can be executed. If \a event_wait_list is NULL, then
 *  this particular command does not wait on any event to complete. If
 *  \a event_wait_list is NULL, \a num_events_in_wait_list must be 0. If
 *  \a event_wait_list is not NULL, the list of events pointed to by
 *  \a event_wait_list must be valid and \a num_events_in_wait_list must be
 *  greater than 0. The events specified in \a event_wait_list act as
 *  synchronization points. The context associated with events in
 *  \a event_wait_list and \a command_queue must be the same. The memory
 *  associated with \a event_wait_list can be reused or freed after the function
 *  returns.
 *
 *  \param num_events_in_wait_list specifies the number of elements in
 *  \a even_wait_list
 *
 *  \param event returns an event object that identifies this particular command
 *  and can be used to query or queue a wait for this particular command to
 *  complete. \a event can be NULL in which case it will not be possible for the
 *  application to query the status of this command or queue a wait for this
 *  command to complete. clEnqueueBarrierWithWaitList can be used instead. If
 *  the \a event_wait_list and the \a event arguments are not NULL, the \a event
 *  argument should not refer to an element of the \a event_wait_list array.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid host
 *    command-queue
 *  - CL_INVALID_CONTEXT if context associated with \a command_queue and
 *    events in \a event_wait_list are not the same
 *  - CL_INVALID_VALUE if \a svm_ptr is NULL.
 *  - CL_INVALID_VALUE if \a size is 0 or if values specified in \a map_flags
 *    are not valid.
 *  - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
 *    \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
 *    \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
 *    are not valid events.
 *  - CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST if the operation is
 *    blocking and the execution status of any of the events in
 *    \a event_wait_list is a negative integer value.
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clEnqueueSVMMap,
              (cl_command_queue command_queue, cl_bool blocking_map, cl_map_flags map_flags,
               void* svm_ptr, size_t size, cl_uint num_events_in_wait_list,
               const cl_event* event_wait_list, cl_event* event)) {
  if (!is_valid(command_queue)) {
    return CL_INVALID_COMMAND_QUEUE;
  }

  if (svm_ptr == NULL) {
    return CL_INVALID_VALUE;
  }

  if (size == 0) {
    return CL_INVALID_VALUE;
  }

  if (!validateMapFlags(map_flags)) {
    return CL_INVALID_VALUE;
  }

  amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
  if (NULL == queue) {
    return CL_INVALID_COMMAND_QUEUE;
  }
  amd::HostQueue& hostQueue = *queue;
  size_t offset = 0;
  amd::Memory* svmMem = NULL;
  if ((queue->device()).isFineGrainedSystem()) {
    // leave blank on purpose for FGS no op
  } else {
    svmMem = amd::MemObjMap::FindMemObj(svm_ptr);
    if (NULL != svmMem) {
      // make sure the context is the same as the context of creation of svm space
      if (hostQueue.context() != svmMem->getContext()) {
        LogWarning("different contexts");
        return CL_INVALID_CONTEXT;
      }

      offset = static_cast<address>(svm_ptr) - static_cast<address>(svmMem->getSvmPtr());
      if (offset < 0 || offset + size > svmMem->getSize()) {
        LogWarning("wrong svm address ");
        return CL_INVALID_VALUE;
      }
      amd::Buffer* srcBuffer = svmMem->asBuffer();

      amd::Coord3D srcSize(size);
      amd::Coord3D srcOffset(offset);
      if (NULL != srcBuffer) {
        if (!srcBuffer->validateRegion(srcOffset, srcSize)) {
          return CL_INVALID_VALUE;
        }
      }

      // Make sure we have memory for the command execution
      device::Memory* mem = svmMem->getDeviceMemory(queue->device());
      if (NULL == mem) {
        LogPrintfError("Can't allocate memory size - 0x%08X bytes!", svmMem->getSize());
        return CL_MEM_OBJECT_ALLOCATION_FAILURE;
      }
      // Attempt to allocate the map target now (whether blocking or non-blocking)
      void* mapPtr = mem->allocMapTarget(srcOffset, srcSize, map_flags);
      if (NULL == mapPtr || mapPtr != svm_ptr) {
        return CL_OUT_OF_RESOURCES;
      }
    }
  }

  amd::Command::EventWaitList eventWaitList;
  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
                                       event_wait_list);
  if (err != CL_SUCCESS) {
    return err;
  }

  amd::Command* command = new amd::SvmMapMemoryCommand(hostQueue, eventWaitList, svmMem, size,
                                                       offset, map_flags, svm_ptr);
  if (command == NULL) {
    return CL_OUT_OF_HOST_MEMORY;
  }
  command->enqueue();

  if (blocking_map) {
    command->awaitCompletion();
  }

  *not_null(event) = as_cl(&command->event());
  if (event == NULL) {
    command->release();
  }
  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief enqueues a command to indicate that the host has completed updating
 * a memory region which was specified in a previous call to clEnqueueSVMUnmap.
 *
 *  \param command_queue is a valid host command-queue.
 *
 *  \param svm_ptr is a pointer that was specified in a previous call to
 *  clEnqueueSVMMap. If \a svm_ptr is allocated using clSVMAlloc then it must be
 *  allocated from the same context from which \a command_queue was created.
 *  Otherwise the behavior is undefined.
 *
 *  \param even_wait_list specifies the events that need to complete before
 *  this particular command can be executed. If \a event_wait_list is NULL, then
 *  this particular command does not wait on any event to complete. If
 *  \a event_wait_list is NULL, \a num_events_in_wait_list must be 0. If
 *  \a event_wait_list is not NULL, the list of events pointed to by
 *  \a event_wait_list must be valid and \a num_events_in_wait_list must be
 *  greater than 0. The events specified in \a event_wait_list act as
 *  synchronization points. The context associated with events in
 *  \a event_wait_list and \a command_queue must be the same. The memory
 *  associated with \a event_wait_list can be reused or freed after the function
 *  returns.
 *
 *  \param num_events_in_wait_list specifies the number of elements in
 *  \a even_wait_list
 *
 *  \param event returns an event object that identifies this particular command
 *  and can be used to query or queue a wait for this particular command to
 *  complete. \a event can be NULL in which case it will not be possible for the
 *  application to query the status of this command or queue a wait for this
 *  command to complete. clEnqueueBarrierWithWaitList can be used instead. If
 *  the \a event_wait_list and the \a event arguments are not NULL, the \a event
 *  argument should not refer to an element of the \a event_wait_list array.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid host
 *    command-queue
 *  - CL_INVALID_CONTEXT if context associated with \a command_queue and
 *    events in \a event_wait_list are not the same
 *  - CL_INVALID_VALUE if \a svm_ptr is NULL.
 *  - CL_INVALID_EVENT_WAIT_LIST if \a event_wait_list is NULL and
 *    \a num_events_in_wait_list > 0, or \a event_wait_list is not NULL and
 *    \a num_events_in_wait_list is 0, or if event objects in \a event_wait_list
 *    are not valid events.
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clEnqueueSVMUnmap,
              (cl_command_queue command_queue, void* svm_ptr, cl_uint num_events_in_wait_list,
               const cl_event* event_wait_list, cl_event* event)) {
  if (!is_valid(command_queue)) {
    return CL_INVALID_COMMAND_QUEUE;
  }

  if (svm_ptr == NULL) {
    return CL_INVALID_VALUE;
  }

  amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
  if (NULL == queue) {
    return CL_INVALID_COMMAND_QUEUE;
  }
  amd::HostQueue& hostQueue = *queue;
  amd::Memory* svmMem = NULL;
  if (!(queue->device()).isFineGrainedSystem()) {
    // check if the ptr is in the svm space
    svmMem = amd::MemObjMap::FindMemObj(svm_ptr);
    // Make sure we have memory for the command execution
    if (NULL != svmMem) {
      // Make sure we have memory for the command execution
      device::Memory* mem = svmMem->getDeviceMemory(queue->device());
      if (NULL == mem) {
        LogPrintfError("Can't allocate memory size - 0x%08X bytes!", svmMem->getSize());
        return CL_INVALID_VALUE;
      }
    }
  }

  amd::Command::EventWaitList eventWaitList;
  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
                                       event_wait_list);
  if (err != CL_SUCCESS) {
    return err;
  }

  amd::Command* command = new amd::SvmUnmapMemoryCommand(hostQueue, eventWaitList, svmMem, svm_ptr);
  if (command == NULL) {
    return CL_OUT_OF_HOST_MEMORY;
  }
  command->enqueue();

  *not_null(event) = as_cl(&command->event());
  if (event == NULL) {
    command->release();
  }
  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief Set the argument value for a specific argument of a kernel to be
 *  a SVM pointer.
 *
 *  \param kernel is a valid kernel object.
 *
 *  \param arg_index is the argument index. Arguments to the kernel are referred
 *  by indices that go from 0 for the leftmost argument to n - 1, where n is the
 *  total number of arguments declared by a kernel.
 *
 *  \param arg_value is the SVM pointer that should be used as the argument
 *  value for argument specified by \a arg_index. The SVM pointer pointed to by
 *  \a arg_value is copied and the \a arg_value pointer can therefore be reused
 *  by the application after clSetKernelArgSVMPointer returns. The SVM pointer
 *  specified is the value used by all API calls that enqueue kernel
 *  (clEnqueueNDRangeKernel) until the argument value is changed by a call to
 *  clSetKernelArgSVMPointer for \a kernel. The SVM pointer can only be used for
 *  arguments that are declared to be a pointer to global or constant memory.
 *  The SVM pointer value must be aligned according to the argument?s type. For
 *  example, if the argument is declared to be global float4 *p, the SVM pointer
 *  value passed for p must be at a minimum aligned to a float4. The SVM pointer
 *  value specified as the argument value can be the pointer returned by
 *  clSVMAlloc or can be a pointer + offset into the SVM region.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_KERNEL if \a kernel is not a valid kernel object
 *  - CL_INVALID_ARG_INDEX if \a arg_index is not a valid argument index
 *  - CL_INVALID_ARG_VALUE if \a arg_value is not a valid value
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clSetKernelArgSVMPointer,
              (cl_kernel kernel, cl_uint arg_index, const void* arg_value)) {
  if (!is_valid(kernel)) {
    return CL_INVALID_KERNEL;
  }

  const amd::KernelSignature& signature = as_amd(kernel)->signature();
  if (arg_index >= signature.numParameters()) {
    return CL_INVALID_ARG_INDEX;
  }

  const amd::KernelParameterDescriptor& desc = signature.at(arg_index);
  if (desc.type_ != T_POINTER ||
      !(desc.addressQualifier_ & (CL_KERNEL_ARG_ADDRESS_GLOBAL | CL_KERNEL_ARG_ADDRESS_CONSTANT))) {
    as_amd(kernel)->parameters().reset(static_cast<size_t>(arg_index));
    return CL_INVALID_ARG_VALUE;
  }

  //! @todo We need to check that the alignment of \a arg_value. For instance,
  // if the argument is of type 'global float4*', then \a arg_value must be
  // aligned to sizeof(float4*). Note that desc.size_ contains the size of the
  // pointer type itself and the size of the pointed type.


  // We do not perform additional pointer validations:
  // -verifying pointers returned by SVMAlloc would imply keeping track
  //  of every allocation range and then matching the pointer against that
  //  range. Note that even if the pointer would look correct, nothing
  //  prevents the user from using an offset within the kernel that would
  //  result on an invalid access.
  // -verifying system pointers (if supported) requires matching the pointer
  //  against the address space of the current process.

  as_amd(kernel)->parameters().set(static_cast<size_t>(arg_index), sizeof(arg_value), &arg_value,
                                   true);
  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief Pass additional information other than argument values to a kernel.
 *
 *  \param kernel is a valid kernel object.
 *
 *  \param param_name specifies the information to be passed to \a kernel. It
 *  must be a cl_kernel_exec_info value.
 *
 *  \param param_value_size specifies the size in bytes of the memory pointed to
 *  by \a param_value.
 *
 *  \param param_value is a pointer to memory where the appropiate values
 *  determined by \a param_name are specified.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function was executed successfully
 *  - CL_INVALID_KERNEL if \a kernel is not a valid kernel object.
 *  - CL_INVALID_VALUE if \a param_name is not valid, if  \a param_value is
 *    NULL or if the size specified by \a param_value_size is not valid
 *  - CL_INVALID_OPERATION if \a param_name is
 *    CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM and \a param_value = CL_TRUE
 *    but no devices in context associated with \a kernel support fine-grained
 *    system SVM allocations
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.0r15
 */
RUNTIME_ENTRY(cl_int, clSetKernelExecInfo, (cl_kernel kernel, cl_kernel_exec_info param_name,
                                            size_t param_value_size, const void* param_value)) {
  if (!is_valid(kernel)) {
    return CL_INVALID_KERNEL;
  }

  if (param_value == NULL) {
    return CL_INVALID_VALUE;
  }

  const amd::Kernel* amdKernel = as_amd(kernel);

  switch (param_name) {
    case CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM:
      if (param_value_size != sizeof(cl_bool)) {
        return CL_INVALID_VALUE;
      } else {
        const bool flag = *(static_cast<const bool*>(param_value));
        const amd::Context* amdContext = &amdKernel->program().context();
        bool foundFineGrainedSystemDevice = false;
        const std::vector<amd::Device*>& devices = amdContext->devices();
        for (const auto it : devices) {
          if (it->info().svmCapabilities_ & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM) {
            foundFineGrainedSystemDevice = true;
            break;
          }
        }
        if (flag && !foundFineGrainedSystemDevice) {
          return CL_INVALID_OPERATION;
        }
        amdKernel->parameters().setSvmSystemPointersSupport(flag ? FGS_YES : FGS_NO);
      }
      break;
    case CL_KERNEL_EXEC_INFO_SVM_PTRS:
      if (param_value_size == 0 || !amd::isMultipleOf(param_value_size, sizeof(void*))) {
        return CL_INVALID_VALUE;
      } else {
        size_t count = param_value_size / sizeof(void*);
        void* const* execInfoArray = reinterpret_cast<void* const*>(param_value);
        for (size_t i = 0; i < count; i++) {
          if (NULL == execInfoArray[i]) {
            return CL_INVALID_VALUE;
          }
        }
        amdKernel->parameters().addSvmPtr(execInfoArray, count);
      }
      break;
    case CL_KERNEL_EXEC_INFO_NEW_VCOP_AMD:
      if (param_value_size != sizeof(cl_bool)) {
        return CL_INVALID_VALUE;
      } else {
        const bool newVcopFlag = (*(reinterpret_cast<const cl_bool*>(param_value))) ? true : false;
        amdKernel->parameters().setExecNewVcop(newVcopFlag);
      }
      break;
    case CL_KERNEL_EXEC_INFO_PFPA_VCOP_AMD:
      if (param_value_size != sizeof(cl_bool)) {
        return CL_INVALID_VALUE;
      } else {
        const bool pfpaVcopFlag = (*(reinterpret_cast<const cl_bool*>(param_value))) ? true : false;
        amdKernel->parameters().setExecPfpaVcop(pfpaVcopFlag);
      }
      break;
    default:
      return CL_INVALID_VALUE;
  }

  return CL_SUCCESS;
}
RUNTIME_EXIT

/*! \brief Enqueues a command to indicate which device a set of ranges of SVM
 *  allocations should be associated with. Once the event returned by
 *  \a clEnqueueSVMMigrateMem has become CL_COMPLETE, the ranges specified by
 *  svm pointers and sizes have been successfully migrated to the device
 *  associated with command queue.
 *  The user is responsible for managing the event dependencies associated with
 *  this command in order to avoid overlapping access to SVM allocations.
 *  Improperly specified event dependencies passed to clEnqueueSVMMigrateMem
 *  could result in undefined results
 *
 *  \param command_queue is a valid host command queue. The specified set of
 *  allocation ranges will be migrated to the OpenCL device associated with
 *  command_queue.
 *
 *  \param num_svm_pointers is the number of pointers in the specified
 *  svm_pointers array, and the number of sizes in the sizes array, if sizes
 *  is not NULL.
 *
 *  \param svm_pointers is a pointer to an array of pointers. Each pointer in
 *  this array must be within an allocation produced by a call to clSVMAlloc.
 *
 *  \param sizes is an array of sizes. The pair svm_pointers[i] and sizes[i]
 *  together define the starting address and number of bytes in a range to be
 *  migrated. sizes may be NULL indicating that every allocation containing
 *  any svm_pointer[i] is to be migrated. Also, if sizes[i] is zero, then the
 *  entire allocation containing svm_pointer[i] is migrated.
 *
 *  \param flags is a bit-field that is used to specify migration options.
 *  Table 5.12 describes the possible values for flags.
 *
 *  \param num_events_in_wait_list specifies the number of event objects in
 *  \a event_wait_list.
 *
 *  \param event_wait_list specifies events that need to complete before this
 *  particular command can be executed. If event_wait_list is NULL, then this
 *  particular command does not wait on any event to complete. If
 *  event_wait_list is NULL, num_events_in_wait_list must be 0. If
 *  event_wait_list is not NULL, the list of events pointed to by
 *  event_wait_list must be valid and num_events_in_wait_list must be greater
 *  than 0. The events specified in event_wait_list act as synchronization
 *  points. The context associated with events in event_wait_list and
 *  command_queue must be the same. The memory associated with
 *  event_wait_list can be reused or freed after the function returns.
 *
 *  \param event an returned event object that identifies this particular write
 *  command and can be used to query or queue a wait for this particular
 *  command to complete. event can be NULL in which case it will not be
 *  possible for the application to query the status of this command or queue
 *  another command that waits for this command to complete. If the
 *  event_wait_list and the event arguments are not NULL, the event argument
 *  should not refer to an element of the event_wait_list array.
 *
 *  \return One of the following values:
 *  - CL_SUCCESS if the function is executed successfully
 *  - CL_INVALID_COMMAND_QUEUE if \a command_queue is not a valid command-queue
 *  - CL_INVALID_VALUE if num_svm_pointers is zero or svm_pointers is NULL
 *  - CL_INVALID_VALUE if sizes[i] is non-zero range [svm_pointers[i],
 *    svm_pointers[i]+sizes[i]) is not contained within an existing clSVMAlloc
 *    allocation
 *  - CL_INVALID_EVENT_WAIT_LIST if event_wait_list is NULL and
 *    num_events_in_wait_list > 0, or event_wait_list is not NULL and
 *    num_events_in_wait_list is 0, or if event objects in event_wait_list are
 *    not valid events
 *  - CL_OUT_OF_RESOURCES if there is a failure to allocate resources required
 *    by the OpenCL implementation on the device.
 *  - CL_OUT_OF_HOST_MEMORY if there is a failure to allocate resources required
 *    by the OpenCL implementation on the host.
 *
 *  \version 2.1r00
 */
RUNTIME_ENTRY(cl_int, clEnqueueSVMMigrateMem,
              (cl_command_queue command_queue, cl_uint num_svm_pointers, const void **svm_pointers,
               const size_t *size, cl_mem_migration_flags flags, cl_uint num_events_in_wait_list,
               const cl_event* event_wait_list, cl_event* event)) {

  if (!is_valid(command_queue)) {
    return CL_INVALID_COMMAND_QUEUE;
  }

  amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
  if (NULL == queue) {
    return CL_INVALID_COMMAND_QUEUE;
  }
  amd::HostQueue& hostQueue = *queue;

  if (num_svm_pointers == 0) {
    LogWarning("invalid parameter \"num_svm_pointers = 0\"");
    return CL_INVALID_VALUE;
  }

  if (svm_pointers == NULL) {
    LogWarning("invalid parameter \"svm_pointers = NULL\"");
    return CL_INVALID_VALUE;
  }

  for (cl_uint i = 0; i < num_svm_pointers; i++) {
    if (svm_pointers[i] == NULL) {
      LogWarning("Null pointers are not allowed");
      return CL_INVALID_VALUE;
    }
  }

  if (flags & ~(CL_MIGRATE_MEM_OBJECT_HOST | CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED)) {
    LogWarning("Invalid flag is specified");
    return CL_INVALID_VALUE;
  }

  std::vector<amd::Memory*> memObjects;
  for (cl_uint i = 0; i < num_svm_pointers; i++) {
    const void* svm_ptr = svm_pointers[i];

    amd::Memory* svmMem = amd::MemObjMap::FindMemObj(svm_ptr);
    if (NULL != svmMem) {
      // make sure the context is the same as the context of creation of svm space
      if (hostQueue.context() != svmMem->getContext()) {
        LogWarning("different contexts");
        return CL_INVALID_CONTEXT;
      }

      // Make sure the specified size[i] is within a valid range
      // TODO: handle the size parameter properly
      size_t svm_size = (size == NULL) ? 0 : size[i];
      size_t offset = reinterpret_cast<const_address>(svm_ptr) - reinterpret_cast<address>(svmMem->getSvmPtr());
      if ((offset + svm_size) > svmMem->getSize()) {
        LogWarning("wrong svm address ");
        return CL_INVALID_VALUE;
      }

      memObjects.push_back(svmMem);
    }
  }

  amd::Command::EventWaitList eventWaitList;
  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
                                       event_wait_list);
  if (err != CL_SUCCESS) {
    return err;
  }

  amd::MigrateMemObjectsCommand* command = new amd::MigrateMemObjectsCommand(
      hostQueue, CL_COMMAND_MIGRATE_MEM_OBJECTS, eventWaitList, memObjects, flags);

  if (command == NULL) {
    return CL_OUT_OF_HOST_MEMORY;
  }

  // Make sure we have memory for the command execution
  if (!command->validateMemory()) {
    delete command;
    return CL_MEM_OBJECT_ALLOCATION_FAILURE;
  }

  command->enqueue();

  *not_null(event) = as_cl(&command->event());
  if (event == NULL) {
    command->release();
  }

  return CL_SUCCESS;
}
RUNTIME_EXIT
/*! @}
 *  @}
 */