diff --git a/dpctl/__init__.py b/dpctl/__init__.py index 409b93ef79..c033a75de2 100644 --- a/dpctl/__init__.py +++ b/dpctl/__init__.py @@ -65,7 +65,7 @@ ) from ._version import get_versions -from .enum_types import backend_type, device_type +from .enum_types import backend_type, device_type, event_status_type __all__ = [ "SyclContext", @@ -88,6 +88,7 @@ ] __all__ += [ "SyclEvent", + "SyclEventRaw", ] __all__ += [ "get_platforms", @@ -112,6 +113,7 @@ __all__ += [ "device_type", "backend_type", + "event_status_type", ] __all__ += [ "get_include", diff --git a/dpctl/_backend.pxd b/dpctl/_backend.pxd index 4b6ca0303b..4af83deb4e 100644 --- a/dpctl/_backend.pxd +++ b/dpctl/_backend.pxd @@ -21,7 +21,7 @@ types defined by dpctl's C API. """ -from libc.stdint cimport int64_t, uint32_t +from libc.stdint cimport int64_t, uint32_t, uint64_t from libcpp cimport bool @@ -104,6 +104,12 @@ cdef extern from "dpctl_sycl_enum_types.h": _L1_cache 'L1_cache', _next_partitionable 'next_partitionable', + ctypedef enum _event_status_type 'DPCTLSyclEventStatusType': + _UNKNOWN_STATUS 'DPCTL_UNKNOWN_STATUS' + _SUBMITTED 'DPCTL_SUBMITTED' + _RUNNING 'DPCTL_RUNNING' + _COMPLETE 'DPCTL_COMPLETE' + cdef extern from "dpctl_sycl_types.h": cdef struct DPCTLOpaqueSyclContext @@ -217,8 +223,25 @@ cdef extern from "dpctl_sycl_device_selector_interface.h": cdef extern from "dpctl_sycl_event_interface.h": + cdef DPCTLSyclEventRef DPCTLEvent_Create() + cdef DPCTLSyclEventRef DPCTLEvent_Copy(const DPCTLSyclEventRef ERef) cdef void DPCTLEvent_Wait(DPCTLSyclEventRef ERef) + cdef void DPCTLEvent_WaitAndThrow(DPCTLSyclEventRef ERef) cdef void DPCTLEvent_Delete(DPCTLSyclEventRef ERef) + cdef _event_status_type DPCTLEvent_GetCommandExecutionStatus(DPCTLSyclEventRef ERef) + cdef _backend_type DPCTLEvent_GetBackend(DPCTLSyclEventRef ERef) + cdef struct DPCTLEventVector + ctypedef DPCTLEventVector *DPCTLEventVectorRef + cdef void DPCTLEventVector_Delete(DPCTLEventVectorRef EVRef) + cdef size_t DPCTLEventVector_Size(DPCTLEventVectorRef EVRef) + cdef DPCTLSyclEventRef DPCTLEventVector_GetAt( + DPCTLEventVectorRef EVRef, + size_t index) + cdef DPCTLEventVectorRef DPCTLEvent_GetWaitList( + DPCTLSyclEventRef ERef) + cdef uint64_t DPCTLEvent_GetProfilingInfoSubmit(DPCTLSyclEventRef ERef) + cdef uint64_t DPCTLEvent_GetProfilingInfoStart(DPCTLSyclEventRef ERef) + cdef uint64_t DPCTLEvent_GetProfilingInfoEnd(DPCTLSyclEventRef ERef) cdef extern from "dpctl_sycl_kernel_interface.h": diff --git a/dpctl/_sycl_event.pxd b/dpctl/_sycl_event.pxd index 64f4b30fac..72d8e4eb3d 100644 --- a/dpctl/_sycl_event.pxd +++ b/dpctl/_sycl_event.pxd @@ -23,13 +23,27 @@ from ._backend cimport DPCTLSyclEventRef -cdef public api class SyclEvent [object PySyclEventObject, type PySyclEventType]: - ''' Wrapper class for a Sycl Event - ''' - cdef DPCTLSyclEventRef _event_ref - cdef list _args +cdef public api class _SyclEvent [ + object Py_SyclEventObject, + type Py_SyclEventType +]: + """ Data owner for SyclEvent + """ + cdef DPCTLSyclEventRef _event_ref + cdef object args + +cdef public api class SyclEvent(_SyclEvent) [ + object PySyclEventObject, + type PySyclEventType +]: + """ Python wrapper class for a ``cl::sycl::event`` + """ @staticmethod - cdef SyclEvent _create (DPCTLSyclEventRef e, list args) - cdef DPCTLSyclEventRef get_event_ref (self) + cdef SyclEvent _create (DPCTLSyclEventRef event, object args=*) + cdef int _init_event_default(self) + cdef int _init_event_from__SyclEvent(self, _SyclEvent other) + cdef int _init_event_from_capsule(self, object caps) + cdef DPCTLSyclEventRef get_event_ref (self) + cdef void _wait (SyclEvent event) cpdef void wait (self) diff --git a/dpctl/_sycl_event.pyx b/dpctl/_sycl_event.pyx index 29b7733913..3e81a188ee 100644 --- a/dpctl/_sycl_event.pyx +++ b/dpctl/_sycl_event.pyx @@ -23,7 +23,32 @@ import logging -from ._backend cimport DPCTLEvent_Delete, DPCTLEvent_Wait, DPCTLSyclEventRef +from cpython cimport pycapsule +from libc.stdint cimport uint64_t +import collections.abc + +from ._backend cimport ( # noqa: E211 + DPCTLEvent_Copy, + DPCTLEvent_Create, + DPCTLEvent_Delete, + DPCTLEvent_GetBackend, + DPCTLEvent_GetCommandExecutionStatus, + DPCTLEvent_GetProfilingInfoEnd, + DPCTLEvent_GetProfilingInfoStart, + DPCTLEvent_GetProfilingInfoSubmit, + DPCTLEvent_GetWaitList, + DPCTLEvent_Wait, + DPCTLEvent_WaitAndThrow, + DPCTLEventVector_Delete, + DPCTLEventVector_GetAt, + DPCTLEventVector_Size, + DPCTLEventVectorRef, + DPCTLSyclEventRef, + _backend_type, + _event_status_type, +) + +from .enum_types import backend_type, event_status_type __all__ = [ "SyclEvent", @@ -39,35 +64,291 @@ cdef api DPCTLSyclEventRef get_event_ref(SyclEvent ev): return ev.get_event_ref() -cdef class SyclEvent: - """ Python wrapper class for cl::sycl::event. - """ +cdef void _event_capsule_deleter(object o): + cdef DPCTLSyclEventRef ERef = NULL + if pycapsule.PyCapsule_IsValid(o, "SyclEventRef"): + ERef = pycapsule.PyCapsule_GetPointer( + o, "SyclEventRef" + ) + DPCTLEvent_Delete(ERef) - @staticmethod - cdef SyclEvent _create(DPCTLSyclEventRef eref, list args): - cdef SyclEvent ret = SyclEvent.__new__(SyclEvent) - ret._event_ref = eref - ret._args = args - return ret + +cdef void _init_helper(_SyclEvent event, DPCTLSyclEventRef ERef): + "Populate attributes of class from opaque reference ERef" + event._event_ref = ERef + + +cdef class _SyclEvent: + """ Data owner for SyclEvent + """ def __dealloc__(self): - self.wait() + DPCTLEvent_Wait(self._event_ref) DPCTLEvent_Delete(self._event_ref) + self.args = None + + +cdef class SyclEvent(_SyclEvent): + """ + SyclEvent(arg=None) + Python class representing ``cl::sycl::event``. There are multiple + ways to create a :class:`dpctl.SyclEvent` object: + + - Invoking the constructor with no arguments creates a ready event + using the default constructor of the ``cl::sycl::event``. + + :Example: + .. code-block:: python + + import dpctl + + # Create a default SyclEvent + e = dpctl.SyclEvent() + + - Invoking the constuctor with a named ``PyCapsule`` with name + **"SyclEventRef"** that carries a pointer to a ``sycl::event`` + object. The capsule will be renamed upon successful consumption + to ensure one-time use. A new named capsule can be constructed by + using :func:`dpctl.SyclEvent._get_capsule` method. + + Args: + arg (optional): Defaults to ``None``. + The argument can be a :class:`dpctl.SyclEvent` + instance, a :class:`dpctl.SyclEvent` instance, or a + named ``PyCapsule`` called **"SyclEventRef"**. + + Raises: + ValueError: If the :class:`dpctl.SyclEvent` object creation failed. + TypeError: In case of incorrect arguments given to constructors, + unexpected types of input arguments, or in the case the input + capsule contained a null pointer or could not be renamed. + """ + + @staticmethod + cdef SyclEvent _create(DPCTLSyclEventRef eref, object args=None): + """" + This function calls DPCTLEvent_Delete(eref). + + The user of this function must pass a copy to keep the + eref argument alive. + """ + cdef _SyclEvent ret = _SyclEvent.__new__(_SyclEvent) + _init_helper(ret, eref) + ret.args=args + return SyclEvent(ret) + + cdef int _init_event_default(self): + self._event_ref = DPCTLEvent_Create() + if (self._event_ref is NULL): + return -1 + self.args=None + return 0 + + cdef int _init_event_from__SyclEvent(self, _SyclEvent other): + self._event_ref = DPCTLEvent_Copy(other._event_ref) + if (self._event_ref is NULL): + return -1 + self.args = other.args + return 0 + + cdef int _init_event_from_capsule(self, object cap): + cdef DPCTLSyclEventRef ERef = NULL + cdef DPCTLSyclEventRef ERef_copy = NULL + cdef int ret = 0 + if pycapsule.PyCapsule_IsValid(cap, "SyclEventRef"): + ERef = pycapsule.PyCapsule_GetPointer( + cap, "SyclEventRef" + ) + if (ERef is NULL): + return -2 + ret = pycapsule.PyCapsule_SetName(cap, "used_SyclEventRef") + if (ret): + return -2 + ERef_copy = DPCTLEvent_Copy(ERef) + if (ERef_copy is NULL): + return -3 + self._event_ref = ERef_copy + self.args = None + return 0 + else: + return -128 + + def __cinit__(self, arg=None): + cdef int ret = 0 + if arg is None: + ret = self._init_event_default() + elif type(arg) is _SyclEvent: + ret = self._init_event_from__SyclEvent(<_SyclEvent> arg) + elif pycapsule.PyCapsule_IsValid(arg, "SyclEventRef"): + ret = self._init_event_from_capsule(arg) + else: + raise TypeError( + "Invalid argument." + ) + if (ret < 0): + if (ret == -1): + raise ValueError("Event failed to be created.") + elif (ret == -2): + raise TypeError( + "Input capsule {} contains a null pointer or could not be" + " renamed".format(arg) + ) + elif (ret == -3): + raise ValueError( + "Internal Error: Could not create a copy of a sycl event." + ) + raise ValueError( + "Unrecognized error code ({}) encountered.".format(ret) + ) cdef DPCTLSyclEventRef get_event_ref(self): - """ Returns the DPCTLSyclEventRef pointer for this class. + """ Returns the `DPCTLSyclEventRef` pointer for this class. """ return self._event_ref - cpdef void wait(self): - DPCTLEvent_Wait(self._event_ref) + @staticmethod + cdef void _wait(SyclEvent event): + DPCTLEvent_WaitAndThrow(event._event_ref) + + @staticmethod + def wait_for(event): + """ Waits for a given event or a sequence of events. + """ + if (isinstance(event, collections.abc.Sequence) and + all((isinstance(el, SyclEvent) for el in event))): + for e in event: + SyclEvent._wait(e) + elif isinstance(event, SyclEvent): + SyclEvent._wait(event) + else: + raise TypeError( + "The passed argument is not a SyclEvent type or " + "a sequence of such objects" + ) def addressof_ref(self): - """ Returns the address of the C API DPCTLSyclEventRef pointer as + """ Returns the address of the C API `DPCTLSyclEventRef` pointer as a size_t. Returns: - The address of the DPCTLSyclEventRef object used to create this - SyclEvent cast to a size_t. + The address of the `DPCTLSyclEventRef` object used to create this + `SyclEvent` cast to a size_t. + """ + return self._event_ref + + def _get_capsule(self): + """ + Returns a copy of the underlying ``cl::sycl::event`` pointer as a void + pointer inside a named ``PyCapsule`` that has the name + **SyclEventRef**. The ownership of the pointer inside the capsule is + passed to the caller, and pointer is deleted when the capsule goes out + of scope. + Returns: + :class:`pycapsule`: A capsule object storing a copy of the + ``cl::sycl::event`` pointer belonging to thus + :class:`dpctl.SyclEvent` instance. + Raises: + ValueError: If the ``DPCTLEvent_Copy`` fails to copy the + ``cl::sycl::event`` pointer. """ - return int(self._event_ref) + cdef DPCTLSyclEventRef ERef = NULL + ERef = DPCTLEvent_Copy(self._event_ref) + if (ERef is NULL): + raise ValueError("SyclEvent copy failed.") + return pycapsule.PyCapsule_New( + ERef, + "SyclEventRef", + &_event_capsule_deleter + ) + + @property + def execution_status(self): + """ Returns the event_status_type enum value for this event. + """ + cdef _event_status_type ESTy = DPCTLEvent_GetCommandExecutionStatus( + self._event_ref + ) + if ESTy == _event_status_type._SUBMITTED: + return event_status_type.submitted + elif ESTy == _event_status_type._RUNNING: + return event_status_type.running + elif ESTy == _event_status_type._COMPLETE: + return event_status_type.complete + else: + raise ValueError("Unknown event status.") + + @property + def backend(self): + """Returns the backend_type enum value for the device + associated with this event. + + Returns: + backend_type: The backend for the device. + """ + cdef _backend_type BE = DPCTLEvent_GetBackend(self._event_ref) + if BE == _backend_type._OPENCL: + return backend_type.opencl + elif BE == _backend_type._LEVEL_ZERO: + return backend_type.level_zero + elif BE == _backend_type._HOST: + return backend_type.host + elif BE == _backend_type._CUDA: + return backend_type.cuda + else: + raise ValueError("Unknown backend type.") + + def get_wait_list(self): + """ + Returns the list of :class:`dpctl.SyclEvent` objects that depend + on this event. + """ + cdef DPCTLEventVectorRef EVRef = DPCTLEvent_GetWaitList( + self.get_event_ref() + ) + cdef size_t num_events + cdef size_t i + cdef DPCTLSyclEventRef ERef + if (EVRef is NULL): + raise ValueError("Internal error: NULL event vector encountered") + num_events = DPCTLEventVector_Size(EVRef) + events = [] + for i in range(num_events): + ERef = DPCTLEventVector_GetAt(EVRef, i) + events.append(SyclEvent._create(ERef, args=None)) + DPCTLEventVector_Delete(EVRef) + return events + + def profiling_info_submit(self): + """ + Returns the 64-bit time value in nanoseconds + when ``cl::sycl::command_group`` was submitted to the queue. + """ + cdef uint64_t profiling_info_submit = 0 + profiling_info_submit = DPCTLEvent_GetProfilingInfoSubmit( + self._event_ref + ) + return profiling_info_submit + + @property + def profiling_info_start(self): + """ + Returns the 64-bit time value in nanoseconds + when ``cl::sycl::command_group`` started execution on the device. + """ + cdef uint64_t profiling_info_start = 0 + profiling_info_start = DPCTLEvent_GetProfilingInfoStart(self._event_ref) + return profiling_info_start + + @property + def profiling_info_end(self): + """ + Returns the 64-bit time value in nanoseconds + when ``cl::sycl::command_group`` finished execution on the device. + """ + cdef uint64_t profiling_info_end = 0 + profiling_info_end = DPCTLEvent_GetProfilingInfoEnd(self._event_ref) + return profiling_info_end + + cpdef void wait(self): + "Synchronously wait for completion of this event." + DPCTLEvent_Wait(self._event_ref) diff --git a/dpctl/enum_types.py b/dpctl/enum_types.py index 2c2bd4edca..de11538417 100644 --- a/dpctl/enum_types.py +++ b/dpctl/enum_types.py @@ -25,6 +25,7 @@ __all__ = [ "device_type", "backend_type", + "event_status_type", ] @@ -71,3 +72,11 @@ class backend_type(Enum): host = auto() level_zero = auto() opencl = auto() + + +class event_status_type(Enum): + + unknown_status = auto() + submitted = auto() + running = auto() + complete = auto() diff --git a/dpctl/tests/test_sycl_event.py b/dpctl/tests/test_sycl_event.py new file mode 100644 index 0000000000..289060ec1b --- /dev/null +++ b/dpctl/tests/test_sycl_event.py @@ -0,0 +1,159 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2021 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +""" Defines unit test cases for the SyclEvent class. +""" + +import numpy as np +import pytest + +import dpctl +import dpctl.memory as dpctl_mem +import dpctl.program as dpctl_prog +from dpctl import event_status_type as esty + +from ._helper import has_cpu + + +def produce_event(profiling=False): + oclSrc = " \ + kernel void add(global int* a) { \ + size_t index = get_global_id(0); \ + a[index] = a[index] + 1; \ + }" + if profiling: + q = dpctl.SyclQueue("opencl:cpu", property="enable_profiling") + else: + q = dpctl.SyclQueue("opencl:cpu") + prog = dpctl_prog.create_program_from_source(q, oclSrc) + addKernel = prog.get_sycl_kernel("add") + + bufBytes = 1024 * np.dtype("i").itemsize + abuf = dpctl_mem.MemoryUSMShared(bufBytes, queue=q) + a = np.ndarray((1024), buffer=abuf, dtype="i") + a[:] = np.arange(1024) + args = [] + + args.append(a.base) + r = [1024] + ev = q.submit(addKernel, args, r) + + return ev + + +def test_create_default_event(): + try: + dpctl.SyclEvent() + except ValueError: + pytest.fail("Failed to create a default event") + + +def test_create_event_from_capsule(): + try: + event = dpctl.SyclEvent() + event_capsule = event._get_capsule() + dpctl.SyclEvent(event_capsule) + except ValueError: + pytest.fail("Failed to create an event from capsule") + + +def test_wait_with_event(): + event = dpctl.SyclEvent() + try: + dpctl.SyclEvent.wait_for(event) + except ValueError: + pytest.fail("Failed to wait_for(event)") + event = dpctl.SyclEvent() + try: + event.wait() + except ValueError: + pytest.fail("Failed to wait for the event") + + +def test_wait_with_list(): + event_1 = dpctl.SyclEvent() + event_2 = dpctl.SyclEvent() + try: + dpctl.SyclEvent.wait_for([event_1, event_2]) + except ValueError: + pytest.fail("Failed to wait for events from the list") + + +def test_execution_status(): + event = dpctl.SyclEvent() + try: + event_status = event.execution_status + except ValueError: + pytest.fail("Failed to get an event status") + assert event_status == esty.complete + + +def test_backend(): + try: + dpctl.SyclEvent().backend + except ValueError: + pytest.fail("Failed to get backend from event") + + +@pytest.mark.skip(reason="event::get_wait_list() method returns wrong result") +def test_get_wait_list(): + if has_cpu(): + oclSrc = " \ + kernel void add_k(global float* a) { \ + size_t index = get_global_id(0); \ + a[index] = a[index] + 1; \ + } \ + kernel void sqrt_k(global float* a) { \ + size_t index = get_global_id(0); \ + a[index] = sqrt(a[index]); \ + } \ + kernel void sin_k(global float* a) { \ + size_t index = get_global_id(0); \ + a[index] = sin(a[index]); \ + }" + q = dpctl.SyclQueue("opencl:cpu") + prog = dpctl_prog.create_program_from_source(q, oclSrc) + addKernel = prog.get_sycl_kernel("add_k") + sqrtKernel = prog.get_sycl_kernel("sqrt_k") + sinKernel = prog.get_sycl_kernel("sin_k") + + bufBytes = 1024 * np.dtype("f").itemsize + abuf = dpctl_mem.MemoryUSMShared(bufBytes, queue=q) + a = np.ndarray((1024), buffer=abuf, dtype="f") + a[:] = np.arange(1024) + args = [] + + args.append(a.base) + r = [1024] + ev_1 = q.submit(addKernel, args, r) + ev_2 = q.submit(sqrtKernel, args, r, dEvents=[ev_1]) + ev_3 = q.submit(sinKernel, args, r, dEvents=[ev_2]) + + try: + wait_list = ev_3.get_wait_list() + except ValueError: + pytest.fail("Failed to get a list of waiting events from SyclEvent") + assert len(wait_list) + + +def test_profiling_info(): + if has_cpu(): + event = produce_event(profiling=True) + assert event.profiling_info_submit + assert event.profiling_info_start + assert event.profiling_info_end + else: + pytest.skip("No OpenCL CPU queues available") diff --git a/examples/python/dppy_kernel.py b/examples/python/dppy_kernel.py new file mode 100644 index 0000000000..3384fd5183 --- /dev/null +++ b/examples/python/dppy_kernel.py @@ -0,0 +1,55 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2021 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + + +import numba_dppy +import numpy as np +from sycl_timer import SyclTimer + +import dpctl + + +@numba_dppy.kernel +def dppy_gemm(a, b, c): + i = numba_dppy.get_global_id(0) + j = numba_dppy.get_global_id(1) + if i >= c.shape[0] or j >= c.shape[1]: + return + c[i, j] = 0 + for k in range(c.shape[0]): + c[i, j] += a[i, k] * b[k, j] + + +X = 1024 +Y = 16 +global_size = X, X + +griddim = X, X +blockdim = Y, Y + +a = np.arange(X * X, dtype=np.float32).reshape(X, X) +b = np.array(np.random.random(X * X), dtype=np.float32).reshape(X, X) +c = np.ones_like(a).reshape(X, X) + +q = dpctl.SyclQueue("opencl:gpu", property="enable_profiling") +with dpctl.device_context(q): + timers = SyclTimer(time_scale=1) + with timers(q): + dppy_gemm[griddim, blockdim](a, b, c) + cc = np.dot(a, b) + host_time, device_time = timers.dt() + print("Wall time: ", host_time, "\n", "Device time: ", device_time) + print(np.allclose(c, cc)) diff --git a/examples/python/sycl_timer.py b/examples/python/sycl_timer.py new file mode 100644 index 0000000000..60422b9ebc --- /dev/null +++ b/examples/python/sycl_timer.py @@ -0,0 +1,61 @@ +# Data Parallel Control (dpctl) +# +# Copyright 2020-2021 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + + +import timeit + +import dpctl + + +class SyclTimer: + def __init__(self, host_time=timeit.default_timer, time_scale=1): + self.timer = host_time + self.time_scale = time_scale + + def __call__(self, queue=None): + if isinstance(queue, dpctl.SyclQueue): + if queue.has_enable_profiling: + self.queue = queue + else: + raise ValueError( + "The queue does not contain the enable_profiling property" + ) + else: + raise ValueError( + "The passed queue must be " + ) + return self.__enter__() + + def __enter__(self): + self.event_start = self.queue.submit_barrier() + self.host_start = self.timer() + return self + + def __exit__(self, *args): + self.event_finish = self.queue.submit_barrier() + self.host_finish = self.timer() + + def dt(self): + self.event_start.wait() + self.event_finish.wait() + return ( + (self.host_finish - self.host_start) * self.time_scale, + ( + self.event_finish.profiling_info_start + - self.event_start.profiling_info_end + ) + * (1e-9 * self.time_scale), + )