def run(self):
dep_signal = drvapi.hsa_signal_t(dependent_signal._id)
roc.hsa_amd_memory_async_copy(
host_dst.device_pointer.value, cpu_ctx._agent._id,
local_memory.device_pointer.value, gpu_ctx._agent._id,
kSize, 1, ctypes.byref(dep_signal), completion_signal)
h2l_start.release() # signal h2l to start
def run(self):
h2l_start.acquire() # to wait until l2h thread has started
roc.hsa_amd_memory_async_copy(local_memory.device_pointer.value,
gpu_ctx._agent._id,
host_src.device_pointer.value,
cpu_ctx._agent._id, kSize, 0,
None,
dependent_signal)
def run(self):
dep_signal = drvapi.hsa_signal_t(dependent_signal._id)
roc.hsa_amd_memory_async_copy(host_dst.device_pointer.value,
cpu_ctx._agent._id,
local_memory.device_pointer.value,
gpu_ctx._agent._id, kSize, 1,
ctypes.byref(dep_signal),
completion_signal)
h2l_start.release() # signal h2l to start
def run(self):
h2l_start.acquire() # to wait until l2h thread has started
roc.hsa_amd_memory_async_copy(
local_memory.device_pointer.value, gpu_ctx._agent._id,
host_src.device_pointer.value, cpu_ctx._agent._id, kSize,
0, None, dependent_signal)
