Introduce CUDA backend specification

[AerialMantis]

Introduce a first draft of a CUDA backend specification covering a mapping of the platform, execution and memories models and interoperability with the CUDA API for SYCL implementations targeting a CUDA backend.

The backend specification has been added as an appendix to the SYCL 2020 specification.

A pull request for introducing a test plan to the SYCL CTS for the CUDA backend interoperability defined in this backend specification is here - KhronosGroup/SYCL-CTS#207

[CLAassistant]

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you all sign our Contributor License Agreement before we can accept your contribution.
5 out of 6 committers have signed the CLA.

✅ npmiller
✅ hdelan
✅ JackAKirk
✅ AidanBeltonS
✅ AerialMantis
❌ t4c1
_{You have signed the CLA already but the status is still pending? Let us recheck it.}

[illuhad]

hipSYCL will not be able to follow this specification because it contradicts with the meanings of e.g. queue and context in hipSYCL and how they are mapped to CUDA. I'm not convinced there is need for this spec if only DPC++ is going to implement it.

[AerialMantis]

Thanks for the feedback @illuhad, so this is just a first draft, which is based initially on the DPC++ CUDA backend implementation, as this is what we are most familiar with, however, we appreciate this may not map in the same way in hipSYCL. Our goal is to converge on a specification which all SYCL implementations which provide a CUDA backend can conform to, so we are keen to collaborate with you on how we can incorporate the design of hipSYCL's CUDA backend into this as well.

To that end we'd be keen to understand where it would be difficult for hipSYCL to conform to the specification as we've currently written it, and we can work with you to adapt it, perhaps there are places where we have been too prescriptive, and we may need to be more flexible?

[illuhad]

Hi @AerialMantis, thanks for the clarification regarding your plans, I'm happy to hear you want to converge on a common specification.
I think for kernel code (execution model etc) there is strong overlap, and any potential differences could probably be fixed easily. The differences are mostly on the runtime side. From the quick read of the spec that I had so far, I think there are three main concerns:

1. hipSYCL deliberately decouples sycl::queue and sycl::context from backend objects. For example, it maintains a pool of CUDA streams, and then distributes the work that comes in through all sycl::queue objects across the pool. This has the advantage that it can already benefit from concurrent kernel execution, concurrent memory copies, or overlap of data transfers and compute even if the user code only utilizes a single sycl::queue. Additionally, the performance behavior is independent of the number of SYCL queues, so the user does not need to worry about how to best utilize queues, and can instead focus on expressing the work as part of the SYCL task graph. Similarly, decoupling context allows for an implicit context management that prevents the user from accidentally spawning new backend contexts (e.g. using the default queue constructor). The downside of this approach is that in general there is no well-defined 1:1 mapping between SYCL queue, context and backend objects, which limits interoperability.
2. Similarly, there might not be a direct relation between backend events and sycl::event, because hipSYCL might elide operations early on in the task graph processing as an optimization, in which case the operation becomes "virtual" in hipSYCL terminology, and then the event is not associated with any real backend event. Maybe it might be possible to at least construct a sycl::event on top of a CUDA event, but that's still unclear. Also, there might be other mechanisms in play for synchronization, such as callbacks instead of CUDA events, and whether the SYCL event comes from the CUDA or e.g. the host backend might also be unclear for the user (e.g. a host task might return an event from the CPU backend). Because of all these complexities, hipSYCL currently does not support event interop.
3. While DPC++ relies exclusively on the CUDA driver API, hipSYCL mostly utilizes the CUDA runtime API, (except for a couple of places where the driver API is used). Most runtime types are a direct typedef for a driver type so e.g. interop with driver types might "just work", but it's unclear how well this guaranteed to be the case and we should check to make sure that we don't miss anything.

Of course we could just make all the problematic interop optional to support both hipSYCL and DPC++ designs, but my understanding is that one of main arguments for such a specification would be precisely to make guarantees for users when they wish to interop with the backend.

[AerialMantis]

@illuhad thank you for the detailed response, this is really helpful.

1. hipSYCL deliberately decouples sycl::queue and sycl::context from backend objects. For example, it maintains a pool of CUDA streams, and then distributes the work that comes in through all sycl::queue objects across the pool. This has the advantage that it can already benefit from concurrent kernel execution, concurrent memory copies, or overlap of data transfers and compute even if the user code only utilizes a single sycl::queue. Additionally, the performance behavior is independent of the number of SYCL queues, so the user does not need to worry about how to best utilize queues, and can instead focus on expressing the work as part of the SYCL task graph. Similarly, decoupling context allows for an implicit context management that prevents the user from accidentally spawning new backend contexts (e.g. using the default queue constructor). The downside of this approach is that in general there is no well-defined 1:1 mapping between SYCL queue, context and backend objects, which limits interoperability.

So this is a really good point, and we thought about this quite a bit. Our thinking for mapping of the SYCL queue to CUDA streams was that it should be loosely defined to allow either a 1:1 mapping or a 1:n mapping, to allow for an implementation to have, for example, different streams to overlap kernel functions and data movement or even to provide a separate stream for interop to make it easier to have concurrent execution with work being done via interop. And we believe this should extend to a n:p mapping where there is a fixed pool of CUDA streams regardless of how many SYCL queues are constructed, we can tweak the definition to ensure this case is included.

For the mapping of the SYCL context to CUDA contexts, we chose a mapping of 1:n in order to support the multi-device context use case, where you have a context representing multiple devices. Since CUDA contexts are associated with a single device, it made sense to allow a SYCL context to map to multiple CUDA contexts. However, if you were to have a pool of CUDA streams then it makes sense that this would naturally extend to the CUDA contexts as well.

For the interop API, we wanted to keep the mapping of the SYCL queue separate somewhat from the interop API requirements in that, even though a SYCL queue could map to multiple CUDA streams, the interop API is still only required to provide a single CUDA stream and the implementation can choose which to provide. I believe this should work for hipSYCL. However, for SYCL context I think it's still necessary to return all CUDA contexts as they are unique for each CUDA device, though perhaps this will still work for hipSYCL if you can retrieve the CUDA contexts from the pool corresponding to the devices of the SYCL context?

Where I think this gets tricky is in the interop API for creating SYCL objects; make_queue and make_context if a CUDA backend is implemented using a pool of CUDA streams or CUDA contexts, then it gets difficult to have the objects created only represent the input CUDA streams or CUDA contexts. However, I see a couple of ways this could work:

• Create a new object with it's own CUDA streams or CUDA contexts which are independent of the pools.
• Add the CUDA stream or CUDA context from make_queue or make_context to the pools, but tag them as only useable by the interop constructed SYCL objects.

2. Similarly, there might not be a direct relation between backend events and sycl::event, because hipSYCL might elide operations early on in the task graph processing as an optimization, in which case the operation becomes "virtual" in hipSYCL terminology, and then the event is not associated with any real backend event. Maybe it might be possible to at least construct a sycl::event on top of a CUDA event, but that's still unclear. Also, there might be other mechanisms in play for synchronization, such as callbacks instead of CUDA events, and whether the SYCL event comes from the CUDA or e.g. the host backend might also be unclear for the user (e.g. a host task might return an event from the CPU backend). Because of all these complexities, hipSYCL currently does not support event interop.

This is a good point, and I think most implementations will do this at some point or another, so it definitely needs to be accounted for, and it's difficult to emulate in CUDA as it's not possible to create a user event to represent virtual or composite operations in the runtime. Perhaps we may need to loosen the mapping of the SYCL event to the CUDA event, and say instead that a SYCL event may have an underlying CUDA event associated with it, but it may not. Then perhaps we could require the interop API to be defined, but make it implementation defined when a CUDA event is returned?

Though as you said, it does reduce the usefulness of having interoperability in this case since users wouldn't be able to rely on it always being available, so perhaps it makes less sense to have interop for the SYCL event. There may be other ways to achieve what you would do with a CUDA event, for example since CUDA streams are in order you could simply wait on the stream, though this is not quite the same it may be sufficient. We'll need to think about this some more.

3. While DPC++ relies exclusively on the CUDA driver API, hipSYCL mostly utilizes the CUDA runtime API, (except for a couple of places where the driver API is used). Most runtime types are a direct typedef for a driver type so e.g. interop with driver types might "just work", but it's unclear how well this guaranteed to be the case and we should check to make sure that we don't miss anything.

That's a good point, we'll need to investigate this and make sure that these types are compatible, otherwise we may need to consider a way of the interop API using one or the other of the types, depending on which API the implementation is using. We'll need to look into this some more.

[illuhad]

Thanks for your thoughts @AerialMantis !

For the interop API, we wanted to keep the mapping of the SYCL queue separate somewhat from the interop API requirements in that, even though a SYCL queue could map to multiple CUDA streams, the interop API is still only required to provide a single CUDA stream and the implementation can choose which to provide. I believe this should work for hipSYCL.

I could of course just return some stream from the pool, but I'm not sure if this might not do more harm than good because I don't think it's what users would expect. In that case, there'd be no guarantees that the returned stream synchronizes in any way with previous operations submitted to the SYCL queue (they may have been scheduled to different streams, even for in-order queues) and neither would subsequent operations be guaranteed to synchronize with anything the user submits to the stream via interop.

However, for SYCL context I think it's still necessary to return all CUDA contexts as they are unique for each CUDA device, though perhaps this will still work for hipSYCL if you can retrieve the CUDA contexts from the pool corresponding to the devices of the SYCL context?

That might work. In most workflows, hipSYCL just uses the default context from the CUDA runtime API, which I think can be extracted if the device is known. There are some caveats since hipSYCL also supports multi-device queues, but those are not covered by the SYCL standard anyway.

Where I think this gets tricky is in the interop API for creating SYCL objects; make_queue and make_context if a CUDA backend is implemented using a pool of CUDA streams or CUDA contexts, then it gets difficult to have the objects created only represent the input CUDA streams or CUDA contexts. However, I see a couple of ways this could work:

I'm not sure those functions make much sense for hipSYCL. The hipSYCL interpretation of e.g. a queue is really more of a high-level interface to the SYCL task graph, with some functionality to manage synchronization of groups of tasks (e.g. wait()). To make this more clear I am pondering whether we should introduce the alias using task_group = queue; in hipSYCL. Relating a task group to a stream does not sound very useful semantically, and adding streams at runtime for interop purposes breaks some of the assumptions that the scheduler currently can make. I'll have to think some more about this. For context it might be even more complicated as we couldn't rely on the CUDA runtime API context anymore in the workflows that use it, and this in turn might break other interop use cases that currently "just work" for SYCL+CUDA interop applications that use the runtime API.

This is a good point, and I think most implementations will do this at some point or another, so it definitely needs to be accounted for, and it's difficult to emulate in CUDA as it's not possible to create a user event to represent virtual or composite operations in the runtime. Perhaps we may need to loosen the mapping of the SYCL event to the CUDA event, and say instead that a SYCL event may have an underlying CUDA event associated with it, but it may not. Then perhaps we could require the interop API to be defined, but make it implementation defined when a CUDA event is returned?

Sure, that could be implemented - assuming this is still useful for users then.

I wonder if maybe we should take a step back to see what we want to accomplish with the spec from a user's point of view to make sure that a common spec that targets the least common denominator between DPC++ and hipSYCL is really better for users than two separate implementation backend specs that are specifically built to match the DPC++ and hipSYCL designs?

[gmlueck]

I think it would be good to clarify that the device returned from this API is just a copy of one of the devices in device::get_devices(). For example, I think you could add wording like this (inspired from similar wording we have for the Level Zero interop spec):

The SYCL execution environment for the CUDA backend contains a fixed number of devices that are enumerated via sycl::device::get_devices(). Calling this function does not create a new device. Rather it merely creates a sycl::device object that is a copy of one of the devices from that enumeration.

[AerialMantis]

The wording for this has been updated to reflect that make_device doesn't create a new platform just returns a device object that is a copy of one of the existing devices, so I think this is resolved now.

[gmlueck]

The Level Zero backend has a similar issue about ownership. We addressed this by adding an ownership field to the interop APIs that create SYCL objects, which allows the application to say whether it keeps ownership of the native handle vs. transferring ownership to the implementation. I think it might make sense to use the same technique here. You can see the wording we used in the Level Zero spec at:

https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md#44-level-zero-handles-ownership-and-thread-safety

[npmiller]

Yes, we looked at that, and I do think it makes sense as well, however it does use an extension header, which is fine as the Level Zero is maintained separately, but it's not ideal if we want the CUDA backend spec to be part of the main SYCL specification.

So we decided to go with a simpler approach for now, for the CUDA backend spec, with the idea in mind that we should probably look into maybe codifying what Level Zero is doing directly in the SYCL interop API, or at least a similar system that could be re-used by the non-OpenCL backends to handle ownership.

[gmlueck]

however it does use an extension header, which is fine as the Level Zero is maintained separately, but it's not ideal if we want the CUDA backend spec to be part of the main SYCL specification.

Sorry, I'm not following. What do you mean by "it does use an extension header"? How does this preclude the interop API from being part of the main SYCL specification?

[npmiller]

It doesn't preclude it, we just thought it might be better done separately as a generic change to the interop API that could be used by all backends.

[gmlueck]

I'm not sure it makes sense to try to solve this in a general way in the core spec, since only some backends have this issue of "ownership". (For example, OpenCL doesn't have this issue because objects are reference counted.) For backends that have this issue (e.g. Level Zero, CUDA), I think it makes sense for each backend to define the mechanism for deciding who "owns" an object. Of course, it would be better if all backends had a similar mechanism, though.

If the Level Zero mechanism seems good to others, can we use that for CUDA too? If it has flaws, we could discuss them and potentially adopt a consistent mechanism across both backends.

[npmiller]

I would argue that the OpenCL backend also has this issue of ownership, it's just that it happens to be easily solved by the OpenCL reference counting. And that the Level Zero approach could easily be implemented by the OpenCL backend, where the ownership transfer would just dictate how SYCL calls clRetain* and clRelease* on the OpenCL objects.

I could be completely wrong on that but looking at the interop API with make_* and get_* it seems to me that ownership questions would naturally come up and I get the feeling that maybe it wasn't part of the original interop API simply because it was solved so easily by reference counting for the OpenCL backend.

[gmlueck]

This paragraph seems like a design document or a description of the implementation. I don't think it's relevant for an API spec.

[hdelan]

This is a valid point. We have omitted/rephrased some things in the latest iteration, which we will be posting shortly.

[AerialMantis]

This wording has been removed, so I think this is resolved now.

[gmlueck]

I assume you are talking about the SYCL types here (e.g. sycl::half), correct?

If so, none of this is specific to CUDA. The core SYCL spec already says that kernels using these types are only compatible with a device the has aspect::fp16. I don't see why we need to say this again in the CUDA backend spec.

[hdelan]

This is correct. We have rephrased this section in the next iteration.

[AerialMantis]

This wording has been updated, so I think this is resolved now.

[gmlueck]

Delete this empty section? It seems like you already talked about extensions above.

[hdelan]

This is a placeholder section for later extensions

[AerialMantis]

This wording has been updated, so I think this is resolved now.

[gmlueck]

This first sentence seems confusing to me. Should it read "Each CUDA enabled device ..."? Why is it important to say that, though? It seems like it is just talking about CUDA in general, not anything specific to SYCL interoperability. Maybe this first sentence should just be deleted?

[AerialMantis]

This wording has been updated to address this, so I think this is resolved now.

[gmlueck]

A question came up recently about the return value of get_native() for buffer. What value is returned if buffer is a sub-buffer? Is it a pointer to the original (non-sub) buffer, or is it an offset to the start of the sub-buffer? It would be nice if there was a table describing the behavior of get_native() for each of these objects. That table would be a good place to clarify things like this.

[gmlueck]

Another question about the interop for buffer ... the OpenCL interop has a vector of cl_mem object because there might be multiple OpenCL devices on the system and the buffer might be represented on more than one device. (At least I think that's the reasoning.) Is it correct that there's just a single CUdeviceptr here, or should there be a vector of them?

[AerialMantis]

So we've been discussing application interop for buffer as well and for a few reasons we came to the conclusion that we should not support it and instead encourage buffer interop via interop_handle::get_native_mem:

• As you mentioned, as a buffer can represent multiple devices it would need to return a vector of CUdeviceptr, but because it allocates data lazily, this would also mean we would either need to return a partial list of memory objects depending on the current state of the runtime or enforce allocation within the get_native.
• If the return value has multiple memory objects, we would need some way to reflect which one has the latest copy of the data.
• A buffer can also represent memory on the host and potentially across multiple backends, and representing this would add further complexity to the return type.

In effect we came to the conclusion we would need to replicate much of the inner workings of a buffer in order to make it really useful to users, and the use cases we had could be supported via host task interop.

[illuhad]

This probably cannot be supported by SYCL implementations relying on the CUDA runtime API instead of the driver API such as hipSYCL.

[AerialMantis]

@illuhad you make a good point, thanks for raising this. I think we could resolve this by clarifying the wording such that for some SYCL implementations the mapping of SYCL context to CUDA context is limited to the CUDA primary context, which would limit the usefulness of the feature when the implementation is based on the runtime API, however, would allow us to maintain a generic interface between the two implementation approaches.

So the wording would be clarified such that get_native<context, backend::cuda> is permitted to return the CUDA primary context and make_context<backend::cuda> is permitted to only accept the CUDA primary context, and throw an exception otherwise.

However, in order to allow users to reason about this there would need to be some kind of query which exposes whether the implementation supports interop with non-primary CUDA contexts. Aside from this particular use case, there is another situation where having such a query would be important. Within a host task, when interoperating with the native objects, it's important to know whether the SYCL implementation initializes non-primary contexts, in order to ensure that the native API calls within the host task are compatible.

I see a couple of ways we could do this:

• A. Introduce a query which tells the user specifically whether the SYCL implementation maps all SYCL contexts to the CUDA primary context or whether unique contexts are allocated, this provides the specific information we need in this case, but if perhaps less scalable.
• B. Introduce a broader query which tells the user whether the SYCL implementation is using the driver API or the runtime API, this could be useful for other places where there is a difference between the drive API and the runtime API, though it may not reflect the case where an implementation combines the two APIs.

[illuhad]

Thanks for your thoughts @AerialMantis! I agree that what you are proposing are possible solutions. I have a hunch that in practice it might be easier to work with for users if they know if an implementation is runtime or driver API based rather than make the query more fine-grained and just about context interop.
E.g. if you know that your CUDA backend is runtime API-based, you can make other assumptions about backend interop. In that case, you can know that backend interop is generally just going to work without having to worry about contexts at all, which I think can be very valuable.
I wonder if it might be easier to understand and clearer if the backend specification just had two flavors, or profiles, so that we can expose the behavior of both approaches in a more obvious way.

I suspect that even if an implementation were to utilize both APIs, you could probably figure out a definition to classify them, most likely around how they behave regarding context interop, and whether the current CUDA device from cudaSetDevice() plays a role.
For example, hipSYCL already uses some driver API components for manual management of PTX modules, but still follows the interop behavior of the runtime API-approach.
Do you think there are use cases where such a classification is not possible?

Also, I wonder: Do we believe that there will actually be a use case to invoke make_context with the same context obtained from get_native()? Or will having this query for runtime API backends just tempt users to do things that are unsupported?
My guess is that the appropriate approach to backend interop might be quite different for a runtime and driver API-based implementation. I guess the core question is: Can we have backend interop in client code that works well across both approaches, or would client code need to be structured differently anyway? In the latter case, it might not be useful to stick to a common backend interop interface.

[gmlueck]

I wonder if it might be easier to understand and clearer if the backend specification just had two flavors, or profiles, so that we can expose the behavior of both approaches in a more obvious way.

Just throwing this out there ... it is also an option to have two different CUDA backends. This means:

• There would be two different enumerators in enum backend.
• Each backend would be responsible for defining its own interop specification (if it wanted to support interop).

Having a single backend interop specification is nice, but only if developers can reasonably write interop code that is portable across both CUDA backends. If we think the two CUDA backends are just too different, it might be easier for users if each one had its own separate interop specification.

[AerialMantis]

@illuhad @gmlueck I think you make a good point, I am wary of introducing different profiles of the CUDA backend specification if possible, I think it would be valuable to have a single specification that all implementations can follow, though I recognize this may not be possible if there are cases we can't handle generically.

I've clarified the wording to state that both APIs are supported. The biggest concern there was regarding the context and whether the primary context is used. I've added a note that the CUcontexts that are returned from interop can include the primary context, and in fact, this is now what the DPC++ CUDA backend does. I suspect we may still need a query for whether the implementation is using the CUDA Driver API or CUDA Runtime API, but I'd like to avoid this if possible to avoid having any bifurcation of the specification.

Also, I wonder: Do we believe that there will actually be a use case to invoke make_context with the same context obtained from get_native()? Or will having this query for runtime API backends just tempt users to do things that are unsupported?

This is a good point, I suspect we don't want to support this, it would like cause havoc with the ownership model.

[fraggamuffin]

still need to discuss how to handle context interrupt

[JackAKirk]

@AerialMantis FYI these advise names seem to be different than those used in DPC++: https://github.com/t4c1/llvm-test-suite/blob/1a5a4e9d72cccf290d88f2498c08314d8bfb271e/SYCL/USM/memadvise_cuda.cpp#L39

Do we need to update one or other of the documentation / dpc++ impl?

[AerialMantis]

Good catch, these are different, but I think we may want to update DPC++ to match what we have in here. I have also updated the definitions here to follow the correct extension convention.

[keryell]

Gordon to be back on this in a few weeks.

[fraggamuffin]

a good F2F

[AerialMantis]

@illuhad apologies that it took me so long to get back to this, I've created a significant update to the pull request and actually moved this to a new PR as I ran into issues rebasing the changes.

I could of course just return some stream from the pool, but I'm not sure if this might not do more harm than good because I don't think it's what users would expect. In that case, there'd be no guarantees that the returned stream synchronizes in any way with previous operations submitted to the SYCL queue (they may have been scheduled to different streams, even for in-order queues) and neither would subsequent operations be guaranteed to synchronize with anything the user submits to the stream via interop.

Regarding your comment about which CUDA stream is returned when interoperating with a SYCL queue, and how to correctly synchronize with this, I clarified the wording such that both the application interop and host task interop require synchronization with the SYCL queue prior to retrieving the native CUDA stream or executing the host task. This is quite strict, and perhaps not the most efficient approach, however, it does ensure correctness, and I think we could address this at a later point with a general improvement to the core specification with a feature like your extension for custom operations or a similar extension that we've been working on at Codeplay.

I'm not sure those functions make much sense for hipSYCL. The hipSYCL interpretation of e.g. a queue is really more of a high-level interface to the SYCL task graph, with some functionality to manage synchronization of groups of tasks (e.g. wait()). To make this more clear I am pondering whether we should introduce the alias using task_group = queue; in hipSYCL. Relating a task group to a stream does not sound very useful semantically, and adding streams at runtime for interop purposes breaks some of the assumptions that the scheduler currently can make. I'll have to think some more about this. For context it might be even more complicated as we couldn't rely on the CUDA runtime API context anymore in the workflows that use it, and this in turn might break other interop use cases that currently "just work" for SYCL+CUDA interop applications that use the runtime API.

So I think the way to support this is to have the SYCL queue or context created from a native CUstream or CUcontext independent of the rest of CUstreams or CUcontexts, so if there is a pool, there would be kept separate from those. The high-level scheduling based on dependencies could still apply as normal.

[AerialMantis]

@illuhad @gmlueck @JackAKirk @hdelan FYI I've had to open a new PR for this as I ran into some issues rebasing this, coming back to it after some time. the latest version is now in #420.

[AerialMantis]

Closing this as a new PR has been opened for this; #420.