I want to add sycl backend device in tvm ref to [device_target_interaction](https://github.com/apache/tvm/blob/main/docs/arch/device_target_interactions.rst and sycl
reference sycl
but now the error occurs in FreeDataSpace. There is a problem that how to understand the sentence “These copies are queued to execute on a specific TVMStreamHandle. However, implementations should not assume that CPU buffers remains valid or accessible after the call to CopyDataFromTo completes.”
Can you give some advice about the problem?
Environment
Linux Ubuntu 18.04 TVM version [Hexagon] Asynchronous DMA support
Steps to reproduce
there is my code
sycl_device_api.cc
void* SYCLWorkspace::AllocDataSpace(Device dev, size_t size, size_t alignment,
DLDataType type_hint) {
this->Init();
VLOG(1) << "sycl device allocating " << size << " bytes share memory";
VLOG(1) << "alloc sycl device id is " << dev.device_id << std::endl;
VLOG(1) << "alloc sycl device type is " << dev.device_type << std::endl;
VLOG(1) << "alloc sycl device alignment is " << alignment << std::endl;
// void* ret = sycl::malloc_shared(size, this->devices[dev.device_id], this->context);
void* ret = nullptr;
if(dev.device_type == kDLCPU ){
ret = sycl::aligned_alloc_host(alignment,size,this->contexts[dev.device_id]);
}else if(dev.device_type == kDLSYCL){
ret = sycl::aligned_alloc_device(alignment,size,this->devices[dev.device_id],this->contexts[dev.device_id]);
}else{
std::cerr<<"unknown device type : "<<dev.device_type<<std::endl;
}
if(ret == nullptr)
LOG(ERROR) << "allgn alloc memory failure!"<<std::endl;
VLOG(1) << "alloc sycl device pointer address is " << ret << std::endl;
return ret;
}
void SYCLWorkspace::FreeDataSpace(Device dev, void* ptr) {
SYCL_CALL(this->GetQueue(dev).wait_and_throw());
if(!IsSYCLDevice(dev)){
VLOG(1) << "free not sycl device : "<<dev.device_type;
LOG(WARNING) << "free not sycl device:"<<dev.device_type;
return ;
}else{
VLOG(1) << "free sycl device id is " << dev.device_id << std::endl;
VLOG(1) << "free sycl device type is " << dev.device_type << std::endl;
VLOG(1) << "free sycl device pointer address is " << ptr << std::endl;
}
sycl::queue queue = this->GetQueue(dev);
sycl::free(ptr, queue);
void SYCLWorkspace::CopyDataFromTo(DLTensor* from, DLTensor* to, TVMStreamHandle stream) {
size_t from_size = GetDataSize(*from);
size_t to_size = GetDataSize(*to);
ICHECK_EQ(from_size, to_size);
ICHECK(IsContiguous(*from) && IsContiguous(*to))
<< "CopyDataFromTo only support contiguous array for now";
size_t from_offset = from->byte_offset;
size_t to_offset = to->byte_offset;
VLOG(1) << "from device " << from->device.device_id << " type : "<< from->device.device_type<<std::endl;
VLOG(1) << "to device " << to->device.device_id << " type : "<< to->device.device_type<<std::endl;
const auto* from_data = static_cast<const uint64_t*>(from->data) + from->byte_offset;
auto* to_data = static_cast<uint64_t*>(to->data) + to->byte_offset;
ICHECK(from_size == to_size) << "TVMArrayCopyFromTo: The size must exactly match";
VLOG(1) << "after convert from device data pointer address : " << from_data << std::endl;
VLOG(1) << "after convert to device data pointer address : " << to_data << std::endl;
if (IsSYCLDevice(from->device) && IsSYCLDevice(to->device)){
auto queue = this->GetQueue(to->device);
auto event = queue.memcpy(to_data,from_data,from_size);
SYCL_CALL(event.wait());
}else if (IsSYCLDevice(from->device) && to->device.device_type == kDLCPU){
auto queue = this->GetQueue(from->device);
auto event = queue.memcpy(to_data,from_data,from_size);
SYCL_CALL(event.wait());
}else if (from->device.device_type == kDLCPU && IsSYCLDevice(to->device)){
auto queue = this->GetQueue(to->device);
auto event = queue.memcpy(to_data,from_data,from_size);
SYCL_CALL(event.wait());
}else {
LOG(FATAL) << "Expect copy from/to SYCL or between SYCL";
}
}
there is my print log
tvm_sycl_vlog.log
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:163: alloc sycl device id is 0
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:164: alloc sycl device type is 17
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:165: alloc sycl device alignment is 64
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:178: alloc sycl device pointer address is 0x7f7c3fbdce00
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:286: from device 0 type : 1
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:287: to device 0 type : 17
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:296: after convert from device data pointer address : 0x5ca2080
[08:54:48] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:297: after convert to device data pointer address : 0x7f7c3fbdce00
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:163: alloc sycl device id is 0
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:164: alloc sycl device type is 17
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:165: alloc sycl device alignment is 64
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:178: alloc sycl device pointer address is 0x7f7c3fbdee00
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:218: free sycl device id is 0
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:219: free sycl device type is 17
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:220: free sycl device pointer address is 0x7f7c21000000
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:218: free sycl device id is 0
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:219: free sycl device type is 17
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:220: free sycl device pointer address is 0x7f7c3fbdce00
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:218: free sycl device id is 0
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:219: free sycl device type is 17
[08:54:49] /home/wzy/tvm-sycl/git-tvm-sycl/tvm/src/runtime/sycl/sycl_device_api.cc:220: free sycl device pointer address is 0x7f7c3fbdee00
python: /home/wzy/sycl_workspace/intel-llvm-new/sycl/plugins/cuda/pi_cuda.cpp:4949: pi_result cuda_piextUSMFree(pi_context, void*): Assertion `type == CU_MEMORYTYPE_DEVICE || type == CU_MEMORYTYPE_HOST' failed.