Zero Copy openmp offload (original) (raw)

August 30, 2023, 5:51pm 13

*to and *from are both dynamically allocated arrays. *from is allocated with the llvm_omp_target_host_mem_alloc where as *to is a device array. The original datatypes of to and from is **void so it might be confusing when I wrote *to, *to is an array not a value.

complete mallochost. *ptr should be *from

void ops_device_mallochost(OPS_instance *instance, void** ptr, size_t bytes) {
#if defined(__clang__)
  *ptr = omp_alloc(bytes, llvm_omp_target_host_mem_alloc);
#else
  *ptr = ops_malloc(bytes);
#endif
}

*to should be *ptr:

void ops_device_malloc(OPS_instance *instance, void** ptr, size_t bytes) {
  *ptr = ops_malloc(bytes);
  int device = omp_get_default_device();

  void* device_ptr = omp_target_alloc(bytes, device);
  omp_target_associate_ptr(*ptr, device_ptr, bytes, 0, device);
}

and this is h2d:

void ops_device_memcpy_h2d(OPS_instance *instance, void** to, void **from, size_t size) {
  int host = omp_get_initial_device();
  int device = omp_get_default_device();

  void* device_ptr = omp_get_mapped_ptr(*to, device);
  omp_target_memcpy(device_ptr, *from, size, 0, 0, device, host);
}

The only option I can think of is

#pragma omp target is_device_ptr(*from)
memcpy(*to, *from, size);

will that work?