wangzy0327 opened a new issue, #13527: URL: https://github.com/apache/tvm/issues/13527
Thanks for participating in the TVM community! We use https://discuss.tvm.ai for any general usage questions and discussions. The issue tracker is used for actionable items such as feature proposals discussion, roadmaps, and bug tracking. You are always welcomed to post on the forum first :smile_cat: Issues that are inactive for a period of time may get closed. We adopt this policy so that we won't lose track of actionable issues that may fall at the bottom of the pile. Feel free to reopen a new one if you feel there is an additional problem that needs attention when an old one gets closed. ### Expected behavior 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](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html) reference sycl https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md 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? ### Actual behavior add relative code in ### Environment Linux Ubuntu 18.04 TVM [version](https://github.com/apache/tvm/tree/7f1856d34f03113dc3a7733c010be43446161944) [Hexagon] Asynchronous DMA support ### Steps to reproduce there is my print log <details> <summary>sycl_device_api.cc</summary> ``` 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"; } } ``` </details> <details> <summary>tvm_sycl_vlog.log</summary> ``` [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. ``` </details> ### Triage Please refer to the list of label tags [here](https://github.com/apache/tvm/wiki/Issue-Triage-Labels) to find the relevant tags and add them below in a bullet format (example below). * needs-triage -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
