Fix description of sycl::free
#758
Conversation
The previous wording that memory is deallocated "without waiting for commands operating on it to be completed" is unnecessary, because calling free with such a command in-progress is already defined as undefined behavior.
The previous wording about "not waiting" could have been read as a guarantee that calls to sycl::free are always non-blocking. In practice, calls to sycl::free may need to block until the device can satisfy the deallocation request (depending on implementation details tied to specific backends). Although the new wording cannot be read this way, adding a non-normative note may help developers to understand that they should not rely on the behavior of specific implementations.
|
Thanks from moving to the new format!
What do you mean?
So I'm not really sure buy removing the |
I mean that the specification didn't actually say "
I'm not sure what you mean. You said that existing CUDA backends are probably blocking (and I agree). That's what I meant about the behavior of existing implementations: there are existing implementations of SYCL that implement |
|
Nevermind... I misread; I was thinking before the spec said it was always Which like you said, is totally incorrect. Some implementations are blocking (and making them nonblocking was more or less impossible previously). So,UB is a good clarification. We can discuss if we want that to be Sorry! |
This patch fixes undefined behavior when freeing memory that might still be in use. Using regular zeMemFree (as was done before this change) is unsafe. L0 spec says that for zeMemFree, the application must ensure the device is not referencing memory before it is freed. SYCL sets 'indirect access' flag for every kernel. This means that each kernel can potentially access any memory and hence it is unsafe to free any memory allocation during kernel execution (if that kernel was submitted at some point in time when that allocation was alive). This replaces the 'indirect access tracking' mechanism used in the legacy adapter (see https://github.com/intel/llvm/blob/4c9b19bdd8b4b95b865522d583b6252bda301d98/unified-runtime/source/adapters/level_zero/context.hpp#L153) Related clarification in the SYCL spec: KhronosGroup/SYCL-Docs#758
This PR makes three changes, which I've split into three commits to simplify review.
The description of
sycl::freeis migrated to the new format.An unnecessary statement about "waiting" is removed. This statement was already covered by (and in fact, inconsistent with!)
sycl::freehaving undefined behavior in the case wheresycl::freeis called while a command is in-progress. Whether an implementation waits for in-progress commands or not doesn't matter, because that usage is invalid.I added a non-normative note to highlight that
sycl::freeis not guaranteed to be blocking or non-blocking. The old wording implied thatsycl::freewas always guaranteed to be non-blocking, but this isn't the case for existing implementations.