Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add coarse grained events extension #754

Merged
merged 8 commits into from
Jun 28, 2022
Merged

Conversation

illuhad
Copy link
Collaborator

@illuhad illuhad commented Jun 22, 2022

This PR adds the coarse-grained event extension. Similarly to the DPC++ discard events extension, it addresses the performance impact of having to create an event for every single SYCL operation due to the submit() signature. However, unlike the DPC++ extension, the events returned with this extension remain valid and support the same functionality as regular events.
Instead, the idea is that, when the user is not interested in using the events to construct efficient task graphs, we can just implement the event functionality by synchronizing with the entire backend queue (e.g. CUDA or HIP streams) for backends that are based on in-order queues. Consequently, a coarse-grained event can be implemented without having to construct any backend event. Basically, a coarse-grained event trades worse synchronization performance for better event construction performance/kernel launch latency.

They can either be enabled either only selectively for individual submissions to a queue using a command group property, or for all submissions of a queue using a queue property.

From the documentation:

HIPSYCL_EXT_COARSE_GRAINED_EVENTS

This extension allows to hint to hipSYCL that events associated with command groups can be more coarse-grained and are allowed to synchronize with potentially more operations.
This can allow hipSYCL to trade less synchronization performance for lighter-weight events, and hence lower kernel launch latency. The main benefit are situations where the returned event from submit is not of particular interest, e.g. in in-order queues when no user synchronization with those events are expected.

For example, a coarse grained event for a backend based on in-order queues (e.g. CUDA or HIP) might just synchronize with the entire HIP or CUDA stream - thereby completely eliding the need to create a new backend event.

Coarse-grained events support the same functionality as regular events.

Coarse-grained events can be requested in two ways:

  1. By passing a property to queue which instructs the queue to construct coarse-grained events for all operations that it processes, and
  2. by passing in a property to an individual command group (see HIPSYCL_EXT_CG_PROPERTY_*). In this case, coarse-grained events can be enabled selectively only for some command groups submitted to a queue.

API Reference

namespace sycl::property::queue {
class hipSYCL_coarse_grained_events {};
}

namespace sycl::property::command_group {
class hipSYCL_coarse_grained_events {};
}

CC @al42and @pszi1ard

@illuhad illuhad marked this pull request as ready for review June 28, 2022 02:51
@illuhad illuhad merged commit 55bf4e7 into develop Jun 28, 2022
@illuhad illuhad deleted the feature/coarse-grained-events branch June 28, 2022 12:43
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

1 participant