diff --git a/lib/python3.12/site-packages/prompt_toolkit/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..7a0c4cee9282e41e46c2a646f748f2b892f4ede1
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/__init__.py
@@ -0,0 +1,54 @@
+"""
+prompt_toolkit
+==============
+
+Author: Jonathan Slenders
+
+Description: prompt_toolkit is a Library for building powerful interactive
+ command lines in Python. It can be a replacement for GNU
+ Readline, but it can be much more than that.
+
+See the examples directory to learn about the usage.
+
+Probably, to get started, you might also want to have a look at
+`prompt_toolkit.shortcuts.prompt`.
+"""
+
+from __future__ import annotations
+
+import re
+from importlib import metadata
+
+# note: this is a bit more lax than the actual pep 440 to allow for a/b/rc/dev without a number
+pep440 = re.compile(
+ r"^([1-9]\d*!)?(0|[1-9]\d*)(\.(0|[1-9]\d*))*((a|b|rc)(0|[1-9]\d*)?)?(\.post(0|[1-9]\d*))?(\.dev(0|[1-9]\d*)?)?$",
+ re.UNICODE,
+)
+from .application import Application
+from .formatted_text import ANSI, HTML
+from .shortcuts import PromptSession, choice, print_formatted_text, prompt
+
+# Don't forget to update in `docs/conf.py`!
+__version__ = metadata.version("prompt_toolkit")
+
+assert pep440.match(__version__)
+
+# Version tuple.
+VERSION = tuple(int(v.rstrip("abrc")) for v in __version__.split(".")[:3])
+
+
+__all__ = [
+ # Application.
+ "Application",
+ # Shortcuts.
+ "prompt",
+ "choice",
+ "PromptSession",
+ "print_formatted_text",
+ # Formatted text.
+ "HTML",
+ "ANSI",
+ # Version info.
+ "__version__",
+ "VERSION",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/auto_suggest.py b/lib/python3.12/site-packages/prompt_toolkit/auto_suggest.py
new file mode 100644
index 0000000000000000000000000000000000000000..73213ba690f56f469aea5c9dd0bd2646d4274398
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/auto_suggest.py
@@ -0,0 +1,177 @@
+"""
+`Fish-style `_ like auto-suggestion.
+
+While a user types input in a certain buffer, suggestions are generated
+(asynchronously.) Usually, they are displayed after the input. When the cursor
+presses the right arrow and the cursor is at the end of the input, the
+suggestion will be inserted.
+
+If you want the auto suggestions to be asynchronous (in a background thread),
+because they take too much time, and could potentially block the event loop,
+then wrap the :class:`.AutoSuggest` instance into a
+:class:`.ThreadedAutoSuggest`.
+"""
+
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from typing import TYPE_CHECKING, Callable
+
+from prompt_toolkit.eventloop import run_in_executor_with_context
+
+from .document import Document
+from .filters import Filter, to_filter
+
+if TYPE_CHECKING:
+ from .buffer import Buffer
+
+__all__ = [
+ "Suggestion",
+ "AutoSuggest",
+ "ThreadedAutoSuggest",
+ "DummyAutoSuggest",
+ "AutoSuggestFromHistory",
+ "ConditionalAutoSuggest",
+ "DynamicAutoSuggest",
+]
+
+
+class Suggestion:
+ """
+ Suggestion returned by an auto-suggest algorithm.
+
+ :param text: The suggestion text.
+ """
+
+ def __init__(self, text: str) -> None:
+ self.text = text
+
+ def __repr__(self) -> str:
+ return f"Suggestion({self.text})"
+
+
+class AutoSuggest(metaclass=ABCMeta):
+ """
+ Base class for auto suggestion implementations.
+ """
+
+ @abstractmethod
+ def get_suggestion(self, buffer: Buffer, document: Document) -> Suggestion | None:
+ """
+ Return `None` or a :class:`.Suggestion` instance.
+
+ We receive both :class:`~prompt_toolkit.buffer.Buffer` and
+ :class:`~prompt_toolkit.document.Document`. The reason is that auto
+ suggestions are retrieved asynchronously. (Like completions.) The
+ buffer text could be changed in the meantime, but ``document`` contains
+ the buffer document like it was at the start of the auto suggestion
+ call. So, from here, don't access ``buffer.text``, but use
+ ``document.text`` instead.
+
+ :param buffer: The :class:`~prompt_toolkit.buffer.Buffer` instance.
+ :param document: The :class:`~prompt_toolkit.document.Document` instance.
+ """
+
+ async def get_suggestion_async(
+ self, buff: Buffer, document: Document
+ ) -> Suggestion | None:
+ """
+ Return a :class:`.Future` which is set when the suggestions are ready.
+ This function can be overloaded in order to provide an asynchronous
+ implementation.
+ """
+ return self.get_suggestion(buff, document)
+
+
+class ThreadedAutoSuggest(AutoSuggest):
+ """
+ Wrapper that runs auto suggestions in a thread.
+ (Use this to prevent the user interface from becoming unresponsive if the
+ generation of suggestions takes too much time.)
+ """
+
+ def __init__(self, auto_suggest: AutoSuggest) -> None:
+ self.auto_suggest = auto_suggest
+
+ def get_suggestion(self, buff: Buffer, document: Document) -> Suggestion | None:
+ return self.auto_suggest.get_suggestion(buff, document)
+
+ async def get_suggestion_async(
+ self, buff: Buffer, document: Document
+ ) -> Suggestion | None:
+ """
+ Run the `get_suggestion` function in a thread.
+ """
+
+ def run_get_suggestion_thread() -> Suggestion | None:
+ return self.get_suggestion(buff, document)
+
+ return await run_in_executor_with_context(run_get_suggestion_thread)
+
+
+class DummyAutoSuggest(AutoSuggest):
+ """
+ AutoSuggest class that doesn't return any suggestion.
+ """
+
+ def get_suggestion(self, buffer: Buffer, document: Document) -> Suggestion | None:
+ return None # No suggestion
+
+
+class AutoSuggestFromHistory(AutoSuggest):
+ """
+ Give suggestions based on the lines in the history.
+ """
+
+ def get_suggestion(self, buffer: Buffer, document: Document) -> Suggestion | None:
+ history = buffer.history
+
+ # Consider only the last line for the suggestion.
+ text = document.text.rsplit("\n", 1)[-1]
+
+ # Only create a suggestion when this is not an empty line.
+ if text.strip():
+ # Find first matching line in history.
+ for string in reversed(list(history.get_strings())):
+ for line in reversed(string.splitlines()):
+ if line.startswith(text):
+ return Suggestion(line[len(text) :])
+
+ return None
+
+
+class ConditionalAutoSuggest(AutoSuggest):
+ """
+ Auto suggest that can be turned on and of according to a certain condition.
+ """
+
+ def __init__(self, auto_suggest: AutoSuggest, filter: bool | Filter) -> None:
+ self.auto_suggest = auto_suggest
+ self.filter = to_filter(filter)
+
+ def get_suggestion(self, buffer: Buffer, document: Document) -> Suggestion | None:
+ if self.filter():
+ return self.auto_suggest.get_suggestion(buffer, document)
+
+ return None
+
+
+class DynamicAutoSuggest(AutoSuggest):
+ """
+ Validator class that can dynamically returns any Validator.
+
+ :param get_validator: Callable that returns a :class:`.Validator` instance.
+ """
+
+ def __init__(self, get_auto_suggest: Callable[[], AutoSuggest | None]) -> None:
+ self.get_auto_suggest = get_auto_suggest
+
+ def get_suggestion(self, buff: Buffer, document: Document) -> Suggestion | None:
+ auto_suggest = self.get_auto_suggest() or DummyAutoSuggest()
+ return auto_suggest.get_suggestion(buff, document)
+
+ async def get_suggestion_async(
+ self, buff: Buffer, document: Document
+ ) -> Suggestion | None:
+ auto_suggest = self.get_auto_suggest() or DummyAutoSuggest()
+ return await auto_suggest.get_suggestion_async(buff, document)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/buffer.py b/lib/python3.12/site-packages/prompt_toolkit/buffer.py
new file mode 100644
index 0000000000000000000000000000000000000000..f5847d4ab932f9d5a8a4c8d1134a0d545e93dcfb
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/buffer.py
@@ -0,0 +1,2029 @@
+"""
+Data structures for the Buffer.
+It holds the text, cursor position, history, etc...
+"""
+
+from __future__ import annotations
+
+import asyncio
+import logging
+import os
+import re
+import shlex
+import shutil
+import subprocess
+import tempfile
+from collections import deque
+from enum import Enum
+from functools import wraps
+from typing import Any, Callable, Coroutine, Iterable, TypeVar, cast
+
+from .application.current import get_app
+from .application.run_in_terminal import run_in_terminal
+from .auto_suggest import AutoSuggest, Suggestion
+from .cache import FastDictCache
+from .clipboard import ClipboardData
+from .completion import (
+ CompleteEvent,
+ Completer,
+ Completion,
+ DummyCompleter,
+ get_common_complete_suffix,
+)
+from .document import Document
+from .eventloop import aclosing
+from .filters import FilterOrBool, to_filter
+from .history import History, InMemoryHistory
+from .search import SearchDirection, SearchState
+from .selection import PasteMode, SelectionState, SelectionType
+from .utils import Event, to_str
+from .validation import ValidationError, Validator
+
+__all__ = [
+ "EditReadOnlyBuffer",
+ "Buffer",
+ "CompletionState",
+ "indent",
+ "unindent",
+ "reshape_text",
+]
+
+logger = logging.getLogger(__name__)
+
+
+class EditReadOnlyBuffer(Exception):
+ "Attempt editing of read-only :class:`.Buffer`."
+
+
+class ValidationState(Enum):
+ "The validation state of a buffer. This is set after the validation."
+
+ VALID = "VALID"
+ INVALID = "INVALID"
+ UNKNOWN = "UNKNOWN"
+
+
+class CompletionState:
+ """
+ Immutable class that contains a completion state.
+ """
+
+ def __init__(
+ self,
+ original_document: Document,
+ completions: list[Completion] | None = None,
+ complete_index: int | None = None,
+ ) -> None:
+ #: Document as it was when the completion started.
+ self.original_document = original_document
+
+ #: List of all the current Completion instances which are possible at
+ #: this point.
+ self.completions = completions or []
+
+ #: Position in the `completions` array.
+ #: This can be `None` to indicate "no completion", the original text.
+ self.complete_index = complete_index # Position in the `_completions` array.
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.original_document!r}, <{len(self.completions)!r}> completions, index={self.complete_index!r})"
+
+ def go_to_index(self, index: int | None) -> None:
+ """
+ Create a new :class:`.CompletionState` object with the new index.
+
+ When `index` is `None` deselect the completion.
+ """
+ if self.completions:
+ assert index is None or 0 <= index < len(self.completions)
+ self.complete_index = index
+
+ def new_text_and_position(self) -> tuple[str, int]:
+ """
+ Return (new_text, new_cursor_position) for this completion.
+ """
+ if self.complete_index is None:
+ return self.original_document.text, self.original_document.cursor_position
+ else:
+ original_text_before_cursor = self.original_document.text_before_cursor
+ original_text_after_cursor = self.original_document.text_after_cursor
+
+ c = self.completions[self.complete_index]
+ if c.start_position == 0:
+ before = original_text_before_cursor
+ else:
+ before = original_text_before_cursor[: c.start_position]
+
+ new_text = before + c.text + original_text_after_cursor
+ new_cursor_position = len(before) + len(c.text)
+ return new_text, new_cursor_position
+
+ @property
+ def current_completion(self) -> Completion | None:
+ """
+ Return the current completion, or return `None` when no completion is
+ selected.
+ """
+ if self.complete_index is not None:
+ return self.completions[self.complete_index]
+ return None
+
+
+_QUOTED_WORDS_RE = re.compile(r"""(\s+|".*?"|'.*?')""")
+
+
+class YankNthArgState:
+ """
+ For yank-last-arg/yank-nth-arg: Keep track of where we are in the history.
+ """
+
+ def __init__(
+ self, history_position: int = 0, n: int = -1, previous_inserted_word: str = ""
+ ) -> None:
+ self.history_position = history_position
+ self.previous_inserted_word = previous_inserted_word
+ self.n = n
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(history_position={self.history_position!r}, n={self.n!r}, previous_inserted_word={self.previous_inserted_word!r})"
+
+
+BufferEventHandler = Callable[["Buffer"], None]
+BufferAcceptHandler = Callable[["Buffer"], bool]
+
+
+class Buffer:
+ """
+ The core data structure that holds the text and cursor position of the
+ current input line and implements all text manipulations on top of it. It
+ also implements the history, undo stack and the completion state.
+
+ :param completer: :class:`~prompt_toolkit.completion.Completer` instance.
+ :param history: :class:`~prompt_toolkit.history.History` instance.
+ :param tempfile_suffix: The tempfile suffix (extension) to be used for the
+ "open in editor" function. For a Python REPL, this would be ".py", so
+ that the editor knows the syntax highlighting to use. This can also be
+ a callable that returns a string.
+ :param tempfile: For more advanced tempfile situations where you need
+ control over the subdirectories and filename. For a Git Commit Message,
+ this would be ".git/COMMIT_EDITMSG", so that the editor knows the syntax
+ highlighting to use. This can also be a callable that returns a string.
+ :param name: Name for this buffer. E.g. DEFAULT_BUFFER. This is mostly
+ useful for key bindings where we sometimes prefer to refer to a buffer
+ by their name instead of by reference.
+ :param accept_handler: Called when the buffer input is accepted. (Usually
+ when the user presses `enter`.) The accept handler receives this
+ `Buffer` as input and should return True when the buffer text should be
+ kept instead of calling reset.
+
+ In case of a `PromptSession` for instance, we want to keep the text,
+ because we will exit the application, and only reset it during the next
+ run.
+ :param max_number_of_completions: Never display more than this number of
+ completions, even when the completer can produce more (limited by
+ default to 10k for performance).
+
+ Events:
+
+ :param on_text_changed: When the buffer text changes. (Callable or None.)
+ :param on_text_insert: When new text is inserted. (Callable or None.)
+ :param on_cursor_position_changed: When the cursor moves. (Callable or None.)
+ :param on_completions_changed: When the completions were changed. (Callable or None.)
+ :param on_suggestion_set: When an auto-suggestion text has been set. (Callable or None.)
+
+ Filters:
+
+ :param complete_while_typing: :class:`~prompt_toolkit.filters.Filter`
+ or `bool`. Decide whether or not to do asynchronous autocompleting while
+ typing.
+ :param validate_while_typing: :class:`~prompt_toolkit.filters.Filter`
+ or `bool`. Decide whether or not to do asynchronous validation while
+ typing.
+ :param enable_history_search: :class:`~prompt_toolkit.filters.Filter` or
+ `bool` to indicate when up-arrow partial string matching is enabled. It
+ is advised to not enable this at the same time as
+ `complete_while_typing`, because when there is an autocompletion found,
+ the up arrows usually browse through the completions, rather than
+ through the history.
+ :param read_only: :class:`~prompt_toolkit.filters.Filter`. When True,
+ changes will not be allowed.
+ :param multiline: :class:`~prompt_toolkit.filters.Filter` or `bool`. When
+ not set, pressing `Enter` will call the `accept_handler`. Otherwise,
+ pressing `Esc-Enter` is required.
+ """
+
+ def __init__(
+ self,
+ completer: Completer | None = None,
+ auto_suggest: AutoSuggest | None = None,
+ history: History | None = None,
+ validator: Validator | None = None,
+ tempfile_suffix: str | Callable[[], str] = "",
+ tempfile: str | Callable[[], str] = "",
+ name: str = "",
+ complete_while_typing: FilterOrBool = False,
+ validate_while_typing: FilterOrBool = False,
+ enable_history_search: FilterOrBool = False,
+ document: Document | None = None,
+ accept_handler: BufferAcceptHandler | None = None,
+ read_only: FilterOrBool = False,
+ multiline: FilterOrBool = True,
+ max_number_of_completions: int = 10000,
+ on_text_changed: BufferEventHandler | None = None,
+ on_text_insert: BufferEventHandler | None = None,
+ on_cursor_position_changed: BufferEventHandler | None = None,
+ on_completions_changed: BufferEventHandler | None = None,
+ on_suggestion_set: BufferEventHandler | None = None,
+ ) -> None:
+ # Accept both filters and booleans as input.
+ enable_history_search = to_filter(enable_history_search)
+ complete_while_typing = to_filter(complete_while_typing)
+ validate_while_typing = to_filter(validate_while_typing)
+ read_only = to_filter(read_only)
+ multiline = to_filter(multiline)
+
+ self.completer = completer or DummyCompleter()
+ self.auto_suggest = auto_suggest
+ self.validator = validator
+ self.tempfile_suffix = tempfile_suffix
+ self.tempfile = tempfile
+ self.name = name
+ self.accept_handler = accept_handler
+
+ # Filters. (Usually, used by the key bindings to drive the buffer.)
+ self.complete_while_typing = complete_while_typing
+ self.validate_while_typing = validate_while_typing
+ self.enable_history_search = enable_history_search
+ self.read_only = read_only
+ self.multiline = multiline
+ self.max_number_of_completions = max_number_of_completions
+
+ # Text width. (For wrapping, used by the Vi 'gq' operator.)
+ self.text_width = 0
+
+ #: The command buffer history.
+ # Note that we shouldn't use a lazy 'or' here. bool(history) could be
+ # False when empty.
+ self.history = InMemoryHistory() if history is None else history
+
+ self.__cursor_position = 0
+
+ # Events
+ self.on_text_changed: Event[Buffer] = Event(self, on_text_changed)
+ self.on_text_insert: Event[Buffer] = Event(self, on_text_insert)
+ self.on_cursor_position_changed: Event[Buffer] = Event(
+ self, on_cursor_position_changed
+ )
+ self.on_completions_changed: Event[Buffer] = Event(self, on_completions_changed)
+ self.on_suggestion_set: Event[Buffer] = Event(self, on_suggestion_set)
+
+ # Document cache. (Avoid creating new Document instances.)
+ self._document_cache: FastDictCache[
+ tuple[str, int, SelectionState | None], Document
+ ] = FastDictCache(Document, size=10)
+
+ # Create completer / auto suggestion / validation coroutines.
+ self._async_suggester = self._create_auto_suggest_coroutine()
+ self._async_completer = self._create_completer_coroutine()
+ self._async_validator = self._create_auto_validate_coroutine()
+
+ # Asyncio task for populating the history.
+ self._load_history_task: asyncio.Future[None] | None = None
+
+ # Reset other attributes.
+ self.reset(document=document)
+
+ def __repr__(self) -> str:
+ if len(self.text) < 15:
+ text = self.text
+ else:
+ text = self.text[:12] + "..."
+
+ return f""
+
+ def reset(
+ self, document: Document | None = None, append_to_history: bool = False
+ ) -> None:
+ """
+ :param append_to_history: Append current input to history first.
+ """
+ if append_to_history:
+ self.append_to_history()
+
+ document = document or Document()
+
+ self.__cursor_position = document.cursor_position
+
+ # `ValidationError` instance. (Will be set when the input is wrong.)
+ self.validation_error: ValidationError | None = None
+ self.validation_state: ValidationState | None = ValidationState.UNKNOWN
+
+ # State of the selection.
+ self.selection_state: SelectionState | None = None
+
+ # Multiple cursor mode. (When we press 'I' or 'A' in visual-block mode,
+ # we can insert text on multiple lines at once. This is implemented by
+ # using multiple cursors.)
+ self.multiple_cursor_positions: list[int] = []
+
+ # When doing consecutive up/down movements, prefer to stay at this column.
+ self.preferred_column: int | None = None
+
+ # State of complete browser
+ # For interactive completion through Ctrl-N/Ctrl-P.
+ self.complete_state: CompletionState | None = None
+
+ # State of Emacs yank-nth-arg completion.
+ self.yank_nth_arg_state: YankNthArgState | None = None # for yank-nth-arg.
+
+ # Remember the document that we had *right before* the last paste
+ # operation. This is used for rotating through the kill ring.
+ self.document_before_paste: Document | None = None
+
+ # Current suggestion.
+ self.suggestion: Suggestion | None = None
+
+ # The history search text. (Used for filtering the history when we
+ # browse through it.)
+ self.history_search_text: str | None = None
+
+ # Undo/redo stacks (stack of `(text, cursor_position)`).
+ self._undo_stack: list[tuple[str, int]] = []
+ self._redo_stack: list[tuple[str, int]] = []
+
+ # Cancel history loader. If history loading was still ongoing.
+ # Cancel the `_load_history_task`, so that next repaint of the
+ # `BufferControl` we will repopulate it.
+ if self._load_history_task is not None:
+ self._load_history_task.cancel()
+ self._load_history_task = None
+
+ #: The working lines. Similar to history, except that this can be
+ #: modified. The user can press arrow_up and edit previous entries.
+ #: Ctrl-C should reset this, and copy the whole history back in here.
+ #: Enter should process the current command and append to the real
+ #: history.
+ self._working_lines: deque[str] = deque([document.text])
+ self.__working_index = 0
+
+ def load_history_if_not_yet_loaded(self) -> None:
+ """
+ Create task for populating the buffer history (if not yet done).
+
+ Note::
+
+ This needs to be called from within the event loop of the
+ application, because history loading is async, and we need to be
+ sure the right event loop is active. Therefor, we call this method
+ in the `BufferControl.create_content`.
+
+ There are situations where prompt_toolkit applications are created
+ in one thread, but will later run in a different thread (Ptpython
+ is one example. The REPL runs in a separate thread, in order to
+ prevent interfering with a potential different event loop in the
+ main thread. The REPL UI however is still created in the main
+ thread.) We could decide to not support creating prompt_toolkit
+ objects in one thread and running the application in a different
+ thread, but history loading is the only place where it matters, and
+ this solves it.
+ """
+ if self._load_history_task is None:
+
+ async def load_history() -> None:
+ async for item in self.history.load():
+ self._working_lines.appendleft(item)
+ self.__working_index += 1
+
+ self._load_history_task = get_app().create_background_task(load_history())
+
+ def load_history_done(f: asyncio.Future[None]) -> None:
+ """
+ Handle `load_history` result when either done, cancelled, or
+ when an exception was raised.
+ """
+ try:
+ f.result()
+ except asyncio.CancelledError:
+ # Ignore cancellation. But handle it, so that we don't get
+ # this traceback.
+ pass
+ except GeneratorExit:
+ # Probably not needed, but we had situations where
+ # `GeneratorExit` was raised in `load_history` during
+ # cancellation.
+ pass
+ except BaseException:
+ # Log error if something goes wrong. (We don't have a
+ # caller to which we can propagate this exception.)
+ logger.exception("Loading history failed")
+
+ self._load_history_task.add_done_callback(load_history_done)
+
+ #
+
+ def _set_text(self, value: str) -> bool:
+ """set text at current working_index. Return whether it changed."""
+ working_index = self.working_index
+ working_lines = self._working_lines
+
+ original_value = working_lines[working_index]
+ working_lines[working_index] = value
+
+ # Return True when this text has been changed.
+ if len(value) != len(original_value):
+ # For Python 2, it seems that when two strings have a different
+ # length and one is a prefix of the other, Python still scans
+ # character by character to see whether the strings are different.
+ # (Some benchmarking showed significant differences for big
+ # documents. >100,000 of lines.)
+ return True
+ elif value != original_value:
+ return True
+ return False
+
+ def _set_cursor_position(self, value: int) -> bool:
+ """Set cursor position. Return whether it changed."""
+ original_position = self.__cursor_position
+ self.__cursor_position = max(0, value)
+
+ return self.__cursor_position != original_position
+
+ @property
+ def text(self) -> str:
+ return self._working_lines[self.working_index]
+
+ @text.setter
+ def text(self, value: str) -> None:
+ """
+ Setting text. (When doing this, make sure that the cursor_position is
+ valid for this text. text/cursor_position should be consistent at any time,
+ otherwise set a Document instead.)
+ """
+ # Ensure cursor position remains within the size of the text.
+ if self.cursor_position > len(value):
+ self.cursor_position = len(value)
+
+ # Don't allow editing of read-only buffers.
+ if self.read_only():
+ raise EditReadOnlyBuffer()
+
+ changed = self._set_text(value)
+
+ if changed:
+ self._text_changed()
+
+ # Reset history search text.
+ # (Note that this doesn't need to happen when working_index
+ # changes, which is when we traverse the history. That's why we
+ # don't do this in `self._text_changed`.)
+ self.history_search_text = None
+
+ @property
+ def cursor_position(self) -> int:
+ return self.__cursor_position
+
+ @cursor_position.setter
+ def cursor_position(self, value: int) -> None:
+ """
+ Setting cursor position.
+ """
+ assert isinstance(value, int)
+
+ # Ensure cursor position is within the size of the text.
+ if value > len(self.text):
+ value = len(self.text)
+ if value < 0:
+ value = 0
+
+ changed = self._set_cursor_position(value)
+
+ if changed:
+ self._cursor_position_changed()
+
+ @property
+ def working_index(self) -> int:
+ return self.__working_index
+
+ @working_index.setter
+ def working_index(self, value: int) -> None:
+ if self.__working_index != value:
+ self.__working_index = value
+ # Make sure to reset the cursor position, otherwise we end up in
+ # situations where the cursor position is out of the bounds of the
+ # text.
+ self.cursor_position = 0
+ self._text_changed()
+
+ def _text_changed(self) -> None:
+ # Remove any validation errors and complete state.
+ self.validation_error = None
+ self.validation_state = ValidationState.UNKNOWN
+ self.complete_state = None
+ self.yank_nth_arg_state = None
+ self.document_before_paste = None
+ self.selection_state = None
+ self.suggestion = None
+ self.preferred_column = None
+
+ # fire 'on_text_changed' event.
+ self.on_text_changed.fire()
+
+ # Input validation.
+ # (This happens on all change events, unlike auto completion, also when
+ # deleting text.)
+ if self.validator and self.validate_while_typing():
+ get_app().create_background_task(self._async_validator())
+
+ def _cursor_position_changed(self) -> None:
+ # Remove any complete state.
+ # (Input validation should only be undone when the cursor position
+ # changes.)
+ self.complete_state = None
+ self.yank_nth_arg_state = None
+ self.document_before_paste = None
+
+ # Unset preferred_column. (Will be set after the cursor movement, if
+ # required.)
+ self.preferred_column = None
+
+ # Note that the cursor position can change if we have a selection the
+ # new position of the cursor determines the end of the selection.
+
+ # fire 'on_cursor_position_changed' event.
+ self.on_cursor_position_changed.fire()
+
+ @property
+ def document(self) -> Document:
+ """
+ Return :class:`~prompt_toolkit.document.Document` instance from the
+ current text, cursor position and selection state.
+ """
+ return self._document_cache[
+ self.text, self.cursor_position, self.selection_state
+ ]
+
+ @document.setter
+ def document(self, value: Document) -> None:
+ """
+ Set :class:`~prompt_toolkit.document.Document` instance.
+
+ This will set both the text and cursor position at the same time, but
+ atomically. (Change events will be triggered only after both have been set.)
+ """
+ self.set_document(value)
+
+ def set_document(self, value: Document, bypass_readonly: bool = False) -> None:
+ """
+ Set :class:`~prompt_toolkit.document.Document` instance. Like the
+ ``document`` property, but accept an ``bypass_readonly`` argument.
+
+ :param bypass_readonly: When True, don't raise an
+ :class:`.EditReadOnlyBuffer` exception, even
+ when the buffer is read-only.
+
+ .. warning::
+
+ When this buffer is read-only and `bypass_readonly` was not passed,
+ the `EditReadOnlyBuffer` exception will be caught by the
+ `KeyProcessor` and is silently suppressed. This is important to
+ keep in mind when writing key bindings, because it won't do what
+ you expect, and there won't be a stack trace. Use try/finally
+ around this function if you need some cleanup code.
+ """
+ # Don't allow editing of read-only buffers.
+ if not bypass_readonly and self.read_only():
+ raise EditReadOnlyBuffer()
+
+ # Set text and cursor position first.
+ text_changed = self._set_text(value.text)
+ cursor_position_changed = self._set_cursor_position(value.cursor_position)
+
+ # Now handle change events. (We do this when text/cursor position is
+ # both set and consistent.)
+ if text_changed:
+ self._text_changed()
+ self.history_search_text = None
+
+ if cursor_position_changed:
+ self._cursor_position_changed()
+
+ @property
+ def is_returnable(self) -> bool:
+ """
+ True when there is something handling accept.
+ """
+ return bool(self.accept_handler)
+
+ # End of
+
+ def save_to_undo_stack(self, clear_redo_stack: bool = True) -> None:
+ """
+ Safe current state (input text and cursor position), so that we can
+ restore it by calling undo.
+ """
+ # Safe if the text is different from the text at the top of the stack
+ # is different. If the text is the same, just update the cursor position.
+ if self._undo_stack and self._undo_stack[-1][0] == self.text:
+ self._undo_stack[-1] = (self._undo_stack[-1][0], self.cursor_position)
+ else:
+ self._undo_stack.append((self.text, self.cursor_position))
+
+ # Saving anything to the undo stack, clears the redo stack.
+ if clear_redo_stack:
+ self._redo_stack = []
+
+ def transform_lines(
+ self,
+ line_index_iterator: Iterable[int],
+ transform_callback: Callable[[str], str],
+ ) -> str:
+ """
+ Transforms the text on a range of lines.
+ When the iterator yield an index not in the range of lines that the
+ document contains, it skips them silently.
+
+ To uppercase some lines::
+
+ new_text = transform_lines(range(5,10), lambda text: text.upper())
+
+ :param line_index_iterator: Iterator of line numbers (int)
+ :param transform_callback: callable that takes the original text of a
+ line, and return the new text for this line.
+
+ :returns: The new text.
+ """
+ # Split lines
+ lines = self.text.split("\n")
+
+ # Apply transformation
+ for index in line_index_iterator:
+ try:
+ lines[index] = transform_callback(lines[index])
+ except IndexError:
+ pass
+
+ return "\n".join(lines)
+
+ def transform_current_line(self, transform_callback: Callable[[str], str]) -> None:
+ """
+ Apply the given transformation function to the current line.
+
+ :param transform_callback: callable that takes a string and return a new string.
+ """
+ document = self.document
+ a = document.cursor_position + document.get_start_of_line_position()
+ b = document.cursor_position + document.get_end_of_line_position()
+ self.text = (
+ document.text[:a]
+ + transform_callback(document.text[a:b])
+ + document.text[b:]
+ )
+
+ def transform_region(
+ self, from_: int, to: int, transform_callback: Callable[[str], str]
+ ) -> None:
+ """
+ Transform a part of the input string.
+
+ :param from_: (int) start position.
+ :param to: (int) end position.
+ :param transform_callback: Callable which accepts a string and returns
+ the transformed string.
+ """
+ assert from_ < to
+
+ self.text = "".join(
+ [
+ self.text[:from_]
+ + transform_callback(self.text[from_:to])
+ + self.text[to:]
+ ]
+ )
+
+ def cursor_left(self, count: int = 1) -> None:
+ self.cursor_position += self.document.get_cursor_left_position(count=count)
+
+ def cursor_right(self, count: int = 1) -> None:
+ self.cursor_position += self.document.get_cursor_right_position(count=count)
+
+ def cursor_up(self, count: int = 1) -> None:
+ """(for multiline edit). Move cursor to the previous line."""
+ original_column = self.preferred_column or self.document.cursor_position_col
+ self.cursor_position += self.document.get_cursor_up_position(
+ count=count, preferred_column=original_column
+ )
+
+ # Remember the original column for the next up/down movement.
+ self.preferred_column = original_column
+
+ def cursor_down(self, count: int = 1) -> None:
+ """(for multiline edit). Move cursor to the next line."""
+ original_column = self.preferred_column or self.document.cursor_position_col
+ self.cursor_position += self.document.get_cursor_down_position(
+ count=count, preferred_column=original_column
+ )
+
+ # Remember the original column for the next up/down movement.
+ self.preferred_column = original_column
+
+ def auto_up(
+ self, count: int = 1, go_to_start_of_line_if_history_changes: bool = False
+ ) -> None:
+ """
+ If we're not on the first line (of a multiline input) go a line up,
+ otherwise go back in history. (If nothing is selected.)
+ """
+ if self.complete_state:
+ self.complete_previous(count=count)
+ elif self.document.cursor_position_row > 0:
+ self.cursor_up(count=count)
+ elif not self.selection_state:
+ self.history_backward(count=count)
+
+ # Go to the start of the line?
+ if go_to_start_of_line_if_history_changes:
+ self.cursor_position += self.document.get_start_of_line_position()
+
+ def auto_down(
+ self, count: int = 1, go_to_start_of_line_if_history_changes: bool = False
+ ) -> None:
+ """
+ If we're not on the last line (of a multiline input) go a line down,
+ otherwise go forward in history. (If nothing is selected.)
+ """
+ if self.complete_state:
+ self.complete_next(count=count)
+ elif self.document.cursor_position_row < self.document.line_count - 1:
+ self.cursor_down(count=count)
+ elif not self.selection_state:
+ self.history_forward(count=count)
+
+ # Go to the start of the line?
+ if go_to_start_of_line_if_history_changes:
+ self.cursor_position += self.document.get_start_of_line_position()
+
+ def delete_before_cursor(self, count: int = 1) -> str:
+ """
+ Delete specified number of characters before cursor and return the
+ deleted text.
+ """
+ assert count >= 0
+ deleted = ""
+
+ if self.cursor_position > 0:
+ deleted = self.text[self.cursor_position - count : self.cursor_position]
+
+ new_text = (
+ self.text[: self.cursor_position - count]
+ + self.text[self.cursor_position :]
+ )
+ new_cursor_position = self.cursor_position - len(deleted)
+
+ # Set new Document atomically.
+ self.document = Document(new_text, new_cursor_position)
+
+ return deleted
+
+ def delete(self, count: int = 1) -> str:
+ """
+ Delete specified number of characters and Return the deleted text.
+ """
+ if self.cursor_position < len(self.text):
+ deleted = self.document.text_after_cursor[:count]
+ self.text = (
+ self.text[: self.cursor_position]
+ + self.text[self.cursor_position + len(deleted) :]
+ )
+ return deleted
+ else:
+ return ""
+
+ def join_next_line(self, separator: str = " ") -> None:
+ """
+ Join the next line to the current one by deleting the line ending after
+ the current line.
+ """
+ if not self.document.on_last_line:
+ self.cursor_position += self.document.get_end_of_line_position()
+ self.delete()
+
+ # Remove spaces.
+ self.text = (
+ self.document.text_before_cursor
+ + separator
+ + self.document.text_after_cursor.lstrip(" ")
+ )
+
+ def join_selected_lines(self, separator: str = " ") -> None:
+ """
+ Join the selected lines.
+ """
+ assert self.selection_state
+
+ # Get lines.
+ from_, to = sorted(
+ [self.cursor_position, self.selection_state.original_cursor_position]
+ )
+
+ before = self.text[:from_]
+ lines = self.text[from_:to].splitlines()
+ after = self.text[to:]
+
+ # Replace leading spaces with just one space.
+ lines = [l.lstrip(" ") + separator for l in lines]
+
+ # Set new document.
+ self.document = Document(
+ text=before + "".join(lines) + after,
+ cursor_position=len(before + "".join(lines[:-1])) - 1,
+ )
+
+ def swap_characters_before_cursor(self) -> None:
+ """
+ Swap the last two characters before the cursor.
+ """
+ pos = self.cursor_position
+
+ if pos >= 2:
+ a = self.text[pos - 2]
+ b = self.text[pos - 1]
+
+ self.text = self.text[: pos - 2] + b + a + self.text[pos:]
+
+ def go_to_history(self, index: int) -> None:
+ """
+ Go to this item in the history.
+ """
+ if index < len(self._working_lines):
+ self.working_index = index
+ self.cursor_position = len(self.text)
+
+ def complete_next(self, count: int = 1, disable_wrap_around: bool = False) -> None:
+ """
+ Browse to the next completions.
+ (Does nothing if there are no completion.)
+ """
+ index: int | None
+
+ if self.complete_state:
+ completions_count = len(self.complete_state.completions)
+
+ if self.complete_state.complete_index is None:
+ index = 0
+ elif self.complete_state.complete_index == completions_count - 1:
+ index = None
+
+ if disable_wrap_around:
+ return
+ else:
+ index = min(
+ completions_count - 1, self.complete_state.complete_index + count
+ )
+ self.go_to_completion(index)
+
+ def complete_previous(
+ self, count: int = 1, disable_wrap_around: bool = False
+ ) -> None:
+ """
+ Browse to the previous completions.
+ (Does nothing if there are no completion.)
+ """
+ index: int | None
+
+ if self.complete_state:
+ if self.complete_state.complete_index == 0:
+ index = None
+
+ if disable_wrap_around:
+ return
+ elif self.complete_state.complete_index is None:
+ index = len(self.complete_state.completions) - 1
+ else:
+ index = max(0, self.complete_state.complete_index - count)
+
+ self.go_to_completion(index)
+
+ def cancel_completion(self) -> None:
+ """
+ Cancel completion, go back to the original text.
+ """
+ if self.complete_state:
+ self.go_to_completion(None)
+ self.complete_state = None
+
+ def _set_completions(self, completions: list[Completion]) -> CompletionState:
+ """
+ Start completions. (Generate list of completions and initialize.)
+
+ By default, no completion will be selected.
+ """
+ self.complete_state = CompletionState(
+ original_document=self.document, completions=completions
+ )
+
+ # Trigger event. This should eventually invalidate the layout.
+ self.on_completions_changed.fire()
+
+ return self.complete_state
+
+ def start_history_lines_completion(self) -> None:
+ """
+ Start a completion based on all the other lines in the document and the
+ history.
+ """
+ found_completions: set[str] = set()
+ completions = []
+
+ # For every line of the whole history, find matches with the current line.
+ current_line = self.document.current_line_before_cursor.lstrip()
+
+ for i, string in enumerate(self._working_lines):
+ for j, l in enumerate(string.split("\n")):
+ l = l.strip()
+ if l and l.startswith(current_line):
+ # When a new line has been found.
+ if l not in found_completions:
+ found_completions.add(l)
+
+ # Create completion.
+ if i == self.working_index:
+ display_meta = "Current, line %s" % (j + 1)
+ else:
+ display_meta = f"History {i + 1}, line {j + 1}"
+
+ completions.append(
+ Completion(
+ text=l,
+ start_position=-len(current_line),
+ display_meta=display_meta,
+ )
+ )
+
+ self._set_completions(completions=completions[::-1])
+ self.go_to_completion(0)
+
+ def go_to_completion(self, index: int | None) -> None:
+ """
+ Select a completion from the list of current completions.
+ """
+ assert self.complete_state
+
+ # Set new completion
+ state = self.complete_state
+ state.go_to_index(index)
+
+ # Set text/cursor position
+ new_text, new_cursor_position = state.new_text_and_position()
+ self.document = Document(new_text, new_cursor_position)
+
+ # (changing text/cursor position will unset complete_state.)
+ self.complete_state = state
+
+ def apply_completion(self, completion: Completion) -> None:
+ """
+ Insert a given completion.
+ """
+ # If there was already a completion active, cancel that one.
+ if self.complete_state:
+ self.go_to_completion(None)
+ self.complete_state = None
+
+ # Insert text from the given completion.
+ self.delete_before_cursor(-completion.start_position)
+ self.insert_text(completion.text)
+
+ def _set_history_search(self) -> None:
+ """
+ Set `history_search_text`.
+ (The text before the cursor will be used for filtering the history.)
+ """
+ if self.enable_history_search():
+ if self.history_search_text is None:
+ self.history_search_text = self.document.text_before_cursor
+ else:
+ self.history_search_text = None
+
+ def _history_matches(self, i: int) -> bool:
+ """
+ True when the current entry matches the history search.
+ (when we don't have history search, it's also True.)
+ """
+ return self.history_search_text is None or self._working_lines[i].startswith(
+ self.history_search_text
+ )
+
+ def history_forward(self, count: int = 1) -> None:
+ """
+ Move forwards through the history.
+
+ :param count: Amount of items to move forward.
+ """
+ self._set_history_search()
+
+ # Go forward in history.
+ found_something = False
+
+ for i in range(self.working_index + 1, len(self._working_lines)):
+ if self._history_matches(i):
+ self.working_index = i
+ count -= 1
+ found_something = True
+ if count == 0:
+ break
+
+ # If we found an entry, move cursor to the end of the first line.
+ if found_something:
+ self.cursor_position = 0
+ self.cursor_position += self.document.get_end_of_line_position()
+
+ def history_backward(self, count: int = 1) -> None:
+ """
+ Move backwards through history.
+ """
+ self._set_history_search()
+
+ # Go back in history.
+ found_something = False
+
+ for i in range(self.working_index - 1, -1, -1):
+ if self._history_matches(i):
+ self.working_index = i
+ count -= 1
+ found_something = True
+ if count == 0:
+ break
+
+ # If we move to another entry, move cursor to the end of the line.
+ if found_something:
+ self.cursor_position = len(self.text)
+
+ def yank_nth_arg(self, n: int | None = None, _yank_last_arg: bool = False) -> None:
+ """
+ Pick nth word from previous history entry (depending on current
+ `yank_nth_arg_state`) and insert it at current position. Rotate through
+ history if called repeatedly. If no `n` has been given, take the first
+ argument. (The second word.)
+
+ :param n: (None or int), The index of the word from the previous line
+ to take.
+ """
+ assert n is None or isinstance(n, int)
+ history_strings = self.history.get_strings()
+
+ if not len(history_strings):
+ return
+
+ # Make sure we have a `YankNthArgState`.
+ if self.yank_nth_arg_state is None:
+ state = YankNthArgState(n=-1 if _yank_last_arg else 1)
+ else:
+ state = self.yank_nth_arg_state
+
+ if n is not None:
+ state.n = n
+
+ # Get new history position.
+ new_pos = state.history_position - 1
+ if -new_pos > len(history_strings):
+ new_pos = -1
+
+ # Take argument from line.
+ line = history_strings[new_pos]
+
+ words = [w.strip() for w in _QUOTED_WORDS_RE.split(line)]
+ words = [w for w in words if w]
+ try:
+ word = words[state.n]
+ except IndexError:
+ word = ""
+
+ # Insert new argument.
+ if state.previous_inserted_word:
+ self.delete_before_cursor(len(state.previous_inserted_word))
+ self.insert_text(word)
+
+ # Save state again for next completion. (Note that the 'insert'
+ # operation from above clears `self.yank_nth_arg_state`.)
+ state.previous_inserted_word = word
+ state.history_position = new_pos
+ self.yank_nth_arg_state = state
+
+ def yank_last_arg(self, n: int | None = None) -> None:
+ """
+ Like `yank_nth_arg`, but if no argument has been given, yank the last
+ word by default.
+ """
+ self.yank_nth_arg(n=n, _yank_last_arg=True)
+
+ def start_selection(
+ self, selection_type: SelectionType = SelectionType.CHARACTERS
+ ) -> None:
+ """
+ Take the current cursor position as the start of this selection.
+ """
+ self.selection_state = SelectionState(self.cursor_position, selection_type)
+
+ def copy_selection(self, _cut: bool = False) -> ClipboardData:
+ """
+ Copy selected text and return :class:`.ClipboardData` instance.
+
+ Notice that this doesn't store the copied data on the clipboard yet.
+ You can store it like this:
+
+ .. code:: python
+
+ data = buffer.copy_selection()
+ get_app().clipboard.set_data(data)
+ """
+ new_document, clipboard_data = self.document.cut_selection()
+ if _cut:
+ self.document = new_document
+
+ self.selection_state = None
+ return clipboard_data
+
+ def cut_selection(self) -> ClipboardData:
+ """
+ Delete selected text and return :class:`.ClipboardData` instance.
+ """
+ return self.copy_selection(_cut=True)
+
+ def paste_clipboard_data(
+ self,
+ data: ClipboardData,
+ paste_mode: PasteMode = PasteMode.EMACS,
+ count: int = 1,
+ ) -> None:
+ """
+ Insert the data from the clipboard.
+ """
+ assert isinstance(data, ClipboardData)
+ assert paste_mode in (PasteMode.VI_BEFORE, PasteMode.VI_AFTER, PasteMode.EMACS)
+
+ original_document = self.document
+ self.document = self.document.paste_clipboard_data(
+ data, paste_mode=paste_mode, count=count
+ )
+
+ # Remember original document. This assignment should come at the end,
+ # because assigning to 'document' will erase it.
+ self.document_before_paste = original_document
+
+ def newline(self, copy_margin: bool = True) -> None:
+ """
+ Insert a line ending at the current position.
+ """
+ if copy_margin:
+ self.insert_text("\n" + self.document.leading_whitespace_in_current_line)
+ else:
+ self.insert_text("\n")
+
+ def insert_line_above(self, copy_margin: bool = True) -> None:
+ """
+ Insert a new line above the current one.
+ """
+ if copy_margin:
+ insert = self.document.leading_whitespace_in_current_line + "\n"
+ else:
+ insert = "\n"
+
+ self.cursor_position += self.document.get_start_of_line_position()
+ self.insert_text(insert)
+ self.cursor_position -= 1
+
+ def insert_line_below(self, copy_margin: bool = True) -> None:
+ """
+ Insert a new line below the current one.
+ """
+ if copy_margin:
+ insert = "\n" + self.document.leading_whitespace_in_current_line
+ else:
+ insert = "\n"
+
+ self.cursor_position += self.document.get_end_of_line_position()
+ self.insert_text(insert)
+
+ def insert_text(
+ self,
+ data: str,
+ overwrite: bool = False,
+ move_cursor: bool = True,
+ fire_event: bool = True,
+ ) -> None:
+ """
+ Insert characters at cursor position.
+
+ :param fire_event: Fire `on_text_insert` event. This is mainly used to
+ trigger autocompletion while typing.
+ """
+ # Original text & cursor position.
+ otext = self.text
+ ocpos = self.cursor_position
+
+ # In insert/text mode.
+ if overwrite:
+ # Don't overwrite the newline itself. Just before the line ending,
+ # it should act like insert mode.
+ overwritten_text = otext[ocpos : ocpos + len(data)]
+ if "\n" in overwritten_text:
+ overwritten_text = overwritten_text[: overwritten_text.find("\n")]
+
+ text = otext[:ocpos] + data + otext[ocpos + len(overwritten_text) :]
+ else:
+ text = otext[:ocpos] + data + otext[ocpos:]
+
+ if move_cursor:
+ cpos = self.cursor_position + len(data)
+ else:
+ cpos = self.cursor_position
+
+ # Set new document.
+ # (Set text and cursor position at the same time. Otherwise, setting
+ # the text will fire a change event before the cursor position has been
+ # set. It works better to have this atomic.)
+ self.document = Document(text, cpos)
+
+ # Fire 'on_text_insert' event.
+ if fire_event: # XXX: rename to `start_complete`.
+ self.on_text_insert.fire()
+
+ # Only complete when "complete_while_typing" is enabled.
+ if self.completer and self.complete_while_typing():
+ get_app().create_background_task(self._async_completer())
+
+ # Call auto_suggest.
+ if self.auto_suggest:
+ get_app().create_background_task(self._async_suggester())
+
+ def undo(self) -> None:
+ # Pop from the undo-stack until we find a text that if different from
+ # the current text. (The current logic of `save_to_undo_stack` will
+ # cause that the top of the undo stack is usually the same as the
+ # current text, so in that case we have to pop twice.)
+ while self._undo_stack:
+ text, pos = self._undo_stack.pop()
+
+ if text != self.text:
+ # Push current text to redo stack.
+ self._redo_stack.append((self.text, self.cursor_position))
+
+ # Set new text/cursor_position.
+ self.document = Document(text, cursor_position=pos)
+ break
+
+ def redo(self) -> None:
+ if self._redo_stack:
+ # Copy current state on undo stack.
+ self.save_to_undo_stack(clear_redo_stack=False)
+
+ # Pop state from redo stack.
+ text, pos = self._redo_stack.pop()
+ self.document = Document(text, cursor_position=pos)
+
+ def validate(self, set_cursor: bool = False) -> bool:
+ """
+ Returns `True` if valid.
+
+ :param set_cursor: Set the cursor position, if an error was found.
+ """
+ # Don't call the validator again, if it was already called for the
+ # current input.
+ if self.validation_state != ValidationState.UNKNOWN:
+ return self.validation_state == ValidationState.VALID
+
+ # Call validator.
+ if self.validator:
+ try:
+ self.validator.validate(self.document)
+ except ValidationError as e:
+ # Set cursor position (don't allow invalid values.)
+ if set_cursor:
+ self.cursor_position = min(
+ max(0, e.cursor_position), len(self.text)
+ )
+
+ self.validation_state = ValidationState.INVALID
+ self.validation_error = e
+ return False
+
+ # Handle validation result.
+ self.validation_state = ValidationState.VALID
+ self.validation_error = None
+ return True
+
+ async def _validate_async(self) -> None:
+ """
+ Asynchronous version of `validate()`.
+ This one doesn't set the cursor position.
+
+ We have both variants, because a synchronous version is required.
+ Handling the ENTER key needs to be completely synchronous, otherwise
+ stuff like type-ahead is going to give very weird results. (People
+ could type input while the ENTER key is still processed.)
+
+ An asynchronous version is required if we have `validate_while_typing`
+ enabled.
+ """
+ while True:
+ # Don't call the validator again, if it was already called for the
+ # current input.
+ if self.validation_state != ValidationState.UNKNOWN:
+ return
+
+ # Call validator.
+ error = None
+ document = self.document
+
+ if self.validator:
+ try:
+ await self.validator.validate_async(self.document)
+ except ValidationError as e:
+ error = e
+
+ # If the document changed during the validation, try again.
+ if self.document != document:
+ continue
+
+ # Handle validation result.
+ if error:
+ self.validation_state = ValidationState.INVALID
+ else:
+ self.validation_state = ValidationState.VALID
+
+ self.validation_error = error
+ get_app().invalidate() # Trigger redraw (display error).
+
+ def append_to_history(self) -> None:
+ """
+ Append the current input to the history.
+ """
+ # Save at the tail of the history. (But don't if the last entry the
+ # history is already the same.)
+ if self.text:
+ history_strings = self.history.get_strings()
+ if not len(history_strings) or history_strings[-1] != self.text:
+ self.history.append_string(self.text)
+
+ def _search(
+ self,
+ search_state: SearchState,
+ include_current_position: bool = False,
+ count: int = 1,
+ ) -> tuple[int, int] | None:
+ """
+ Execute search. Return (working_index, cursor_position) tuple when this
+ search is applied. Returns `None` when this text cannot be found.
+ """
+ assert count > 0
+
+ text = search_state.text
+ direction = search_state.direction
+ ignore_case = search_state.ignore_case()
+
+ def search_once(
+ working_index: int, document: Document
+ ) -> tuple[int, Document] | None:
+ """
+ Do search one time.
+ Return (working_index, document) or `None`
+ """
+ if direction == SearchDirection.FORWARD:
+ # Try find at the current input.
+ new_index = document.find(
+ text,
+ include_current_position=include_current_position,
+ ignore_case=ignore_case,
+ )
+
+ if new_index is not None:
+ return (
+ working_index,
+ Document(document.text, document.cursor_position + new_index),
+ )
+ else:
+ # No match, go forward in the history. (Include len+1 to wrap around.)
+ # (Here we should always include all cursor positions, because
+ # it's a different line.)
+ for i in range(working_index + 1, len(self._working_lines) + 1):
+ i %= len(self._working_lines)
+
+ document = Document(self._working_lines[i], 0)
+ new_index = document.find(
+ text, include_current_position=True, ignore_case=ignore_case
+ )
+ if new_index is not None:
+ return (i, Document(document.text, new_index))
+ else:
+ # Try find at the current input.
+ new_index = document.find_backwards(text, ignore_case=ignore_case)
+
+ if new_index is not None:
+ return (
+ working_index,
+ Document(document.text, document.cursor_position + new_index),
+ )
+ else:
+ # No match, go back in the history. (Include -1 to wrap around.)
+ for i in range(working_index - 1, -2, -1):
+ i %= len(self._working_lines)
+
+ document = Document(
+ self._working_lines[i], len(self._working_lines[i])
+ )
+ new_index = document.find_backwards(
+ text, ignore_case=ignore_case
+ )
+ if new_index is not None:
+ return (
+ i,
+ Document(document.text, len(document.text) + new_index),
+ )
+ return None
+
+ # Do 'count' search iterations.
+ working_index = self.working_index
+ document = self.document
+ for _ in range(count):
+ result = search_once(working_index, document)
+ if result is None:
+ return None # Nothing found.
+ else:
+ working_index, document = result
+
+ return (working_index, document.cursor_position)
+
+ def document_for_search(self, search_state: SearchState) -> Document:
+ """
+ Return a :class:`~prompt_toolkit.document.Document` instance that has
+ the text/cursor position for this search, if we would apply it. This
+ will be used in the
+ :class:`~prompt_toolkit.layout.BufferControl` to display feedback while
+ searching.
+ """
+ search_result = self._search(search_state, include_current_position=True)
+
+ if search_result is None:
+ return self.document
+ else:
+ working_index, cursor_position = search_result
+
+ # Keep selection, when `working_index` was not changed.
+ if working_index == self.working_index:
+ selection = self.selection_state
+ else:
+ selection = None
+
+ return Document(
+ self._working_lines[working_index], cursor_position, selection=selection
+ )
+
+ def get_search_position(
+ self,
+ search_state: SearchState,
+ include_current_position: bool = True,
+ count: int = 1,
+ ) -> int:
+ """
+ Get the cursor position for this search.
+ (This operation won't change the `working_index`. It's won't go through
+ the history. Vi text objects can't span multiple items.)
+ """
+ search_result = self._search(
+ search_state, include_current_position=include_current_position, count=count
+ )
+
+ if search_result is None:
+ return self.cursor_position
+ else:
+ working_index, cursor_position = search_result
+ return cursor_position
+
+ def apply_search(
+ self,
+ search_state: SearchState,
+ include_current_position: bool = True,
+ count: int = 1,
+ ) -> None:
+ """
+ Apply search. If something is found, set `working_index` and
+ `cursor_position`.
+ """
+ search_result = self._search(
+ search_state, include_current_position=include_current_position, count=count
+ )
+
+ if search_result is not None:
+ working_index, cursor_position = search_result
+ self.working_index = working_index
+ self.cursor_position = cursor_position
+
+ def exit_selection(self) -> None:
+ self.selection_state = None
+
+ def _editor_simple_tempfile(self) -> tuple[str, Callable[[], None]]:
+ """
+ Simple (file) tempfile implementation.
+ Return (tempfile, cleanup_func).
+ """
+ suffix = to_str(self.tempfile_suffix)
+ descriptor, filename = tempfile.mkstemp(suffix)
+
+ os.write(descriptor, self.text.encode("utf-8"))
+ os.close(descriptor)
+
+ def cleanup() -> None:
+ os.unlink(filename)
+
+ return filename, cleanup
+
+ def _editor_complex_tempfile(self) -> tuple[str, Callable[[], None]]:
+ # Complex (directory) tempfile implementation.
+ headtail = to_str(self.tempfile)
+ if not headtail:
+ # Revert to simple case.
+ return self._editor_simple_tempfile()
+ headtail = str(headtail)
+
+ # Try to make according to tempfile logic.
+ head, tail = os.path.split(headtail)
+ if os.path.isabs(head):
+ head = head[1:]
+
+ dirpath = tempfile.mkdtemp()
+ if head:
+ dirpath = os.path.join(dirpath, head)
+ # Assume there is no issue creating dirs in this temp dir.
+ os.makedirs(dirpath)
+
+ # Open the filename and write current text.
+ filename = os.path.join(dirpath, tail)
+ with open(filename, "w", encoding="utf-8") as fh:
+ fh.write(self.text)
+
+ def cleanup() -> None:
+ shutil.rmtree(dirpath)
+
+ return filename, cleanup
+
+ def open_in_editor(self, validate_and_handle: bool = False) -> asyncio.Task[None]:
+ """
+ Open code in editor.
+
+ This returns a future, and runs in a thread executor.
+ """
+ if self.read_only():
+ raise EditReadOnlyBuffer()
+
+ # Write current text to temporary file
+ if self.tempfile:
+ filename, cleanup_func = self._editor_complex_tempfile()
+ else:
+ filename, cleanup_func = self._editor_simple_tempfile()
+
+ async def run() -> None:
+ try:
+ # Open in editor
+ # (We need to use `run_in_terminal`, because not all editors go to
+ # the alternate screen buffer, and some could influence the cursor
+ # position.)
+ success = await run_in_terminal(
+ lambda: self._open_file_in_editor(filename), in_executor=True
+ )
+
+ # Read content again.
+ if success:
+ with open(filename, "rb") as f:
+ text = f.read().decode("utf-8")
+
+ # Drop trailing newline. (Editors are supposed to add it at the
+ # end, but we don't need it.)
+ if text.endswith("\n"):
+ text = text[:-1]
+
+ self.document = Document(text=text, cursor_position=len(text))
+
+ # Accept the input.
+ if validate_and_handle:
+ self.validate_and_handle()
+
+ finally:
+ # Clean up temp dir/file.
+ cleanup_func()
+
+ return get_app().create_background_task(run())
+
+ def _open_file_in_editor(self, filename: str) -> bool:
+ """
+ Call editor executable.
+
+ Return True when we received a zero return code.
+ """
+ # If the 'VISUAL' or 'EDITOR' environment variable has been set, use that.
+ # Otherwise, fall back to the first available editor that we can find.
+ visual = os.environ.get("VISUAL")
+ editor = os.environ.get("EDITOR")
+
+ editors = [
+ visual,
+ editor,
+ # Order of preference.
+ "/usr/bin/editor",
+ "/usr/bin/nano",
+ "/usr/bin/pico",
+ "/usr/bin/vi",
+ "/usr/bin/emacs",
+ ]
+
+ for e in editors:
+ if e:
+ try:
+ # Use 'shlex.split()', because $VISUAL can contain spaces
+ # and quotes.
+ returncode = subprocess.call(shlex.split(e) + [filename])
+ return returncode == 0
+
+ except OSError:
+ # Executable does not exist, try the next one.
+ pass
+
+ return False
+
+ def start_completion(
+ self,
+ select_first: bool = False,
+ select_last: bool = False,
+ insert_common_part: bool = False,
+ complete_event: CompleteEvent | None = None,
+ ) -> None:
+ """
+ Start asynchronous autocompletion of this buffer.
+ (This will do nothing if a previous completion was still in progress.)
+ """
+ # Only one of these options can be selected.
+ assert select_first + select_last + insert_common_part <= 1
+
+ get_app().create_background_task(
+ self._async_completer(
+ select_first=select_first,
+ select_last=select_last,
+ insert_common_part=insert_common_part,
+ complete_event=complete_event
+ or CompleteEvent(completion_requested=True),
+ )
+ )
+
+ def _create_completer_coroutine(self) -> Callable[..., Coroutine[Any, Any, None]]:
+ """
+ Create function for asynchronous autocompletion.
+
+ (This consumes the asynchronous completer generator, which possibly
+ runs the completion algorithm in another thread.)
+ """
+
+ def completion_does_nothing(document: Document, completion: Completion) -> bool:
+ """
+ Return `True` if applying this completion doesn't have any effect.
+ (When it doesn't insert any new text.
+ """
+ text_before_cursor = document.text_before_cursor
+ replaced_text = text_before_cursor[
+ len(text_before_cursor) + completion.start_position :
+ ]
+ return replaced_text == completion.text
+
+ @_only_one_at_a_time
+ async def async_completer(
+ select_first: bool = False,
+ select_last: bool = False,
+ insert_common_part: bool = False,
+ complete_event: CompleteEvent | None = None,
+ ) -> None:
+ document = self.document
+ complete_event = complete_event or CompleteEvent(text_inserted=True)
+
+ # Don't complete when we already have completions.
+ if self.complete_state or not self.completer:
+ return
+
+ # Create an empty CompletionState.
+ complete_state = CompletionState(original_document=self.document)
+ self.complete_state = complete_state
+
+ def proceed() -> bool:
+ """Keep retrieving completions. Input text has not yet changed
+ while generating completions."""
+ return self.complete_state == complete_state
+
+ refresh_needed = asyncio.Event()
+
+ async def refresh_while_loading() -> None:
+ """Background loop to refresh the UI at most 3 times a second
+ while the completion are loading. Calling
+ `on_completions_changed.fire()` for every completion that we
+ receive is too expensive when there are many completions. (We
+ could tune `Application.max_render_postpone_time` and
+ `Application.min_redraw_interval`, but having this here is a
+ better approach.)
+ """
+ while True:
+ self.on_completions_changed.fire()
+ refresh_needed.clear()
+ await asyncio.sleep(0.3)
+ await refresh_needed.wait()
+
+ refresh_task = asyncio.ensure_future(refresh_while_loading())
+ try:
+ # Load.
+ async with aclosing(
+ self.completer.get_completions_async(document, complete_event)
+ ) as async_generator:
+ async for completion in async_generator:
+ complete_state.completions.append(completion)
+ refresh_needed.set()
+
+ # If the input text changes, abort.
+ if not proceed():
+ break
+
+ # Always stop at 10k completions.
+ if (
+ len(complete_state.completions)
+ >= self.max_number_of_completions
+ ):
+ break
+ finally:
+ refresh_task.cancel()
+
+ # Refresh one final time after we got everything.
+ self.on_completions_changed.fire()
+
+ completions = complete_state.completions
+
+ # When there is only one completion, which has nothing to add, ignore it.
+ if len(completions) == 1 and completion_does_nothing(
+ document, completions[0]
+ ):
+ del completions[:]
+
+ # Set completions if the text was not yet changed.
+ if proceed():
+ # When no completions were found, or when the user selected
+ # already a completion by using the arrow keys, don't do anything.
+ if (
+ not self.complete_state
+ or self.complete_state.complete_index is not None
+ ):
+ return
+
+ # When there are no completions, reset completion state anyway.
+ if not completions:
+ self.complete_state = None
+ # Render the ui if the completion menu was shown
+ # it is needed especially if there is one completion and it was deleted.
+ self.on_completions_changed.fire()
+ return
+
+ # Select first/last or insert common part, depending on the key
+ # binding. (For this we have to wait until all completions are
+ # loaded.)
+
+ if select_first:
+ self.go_to_completion(0)
+
+ elif select_last:
+ self.go_to_completion(len(completions) - 1)
+
+ elif insert_common_part:
+ common_part = get_common_complete_suffix(document, completions)
+ if common_part:
+ # Insert the common part, update completions.
+ self.insert_text(common_part)
+ if len(completions) > 1:
+ # (Don't call `async_completer` again, but
+ # recalculate completions. See:
+ # https://github.com/ipython/ipython/issues/9658)
+ completions[:] = [
+ c.new_completion_from_position(len(common_part))
+ for c in completions
+ ]
+
+ self._set_completions(completions=completions)
+ else:
+ self.complete_state = None
+ else:
+ # When we were asked to insert the "common"
+ # prefix, but there was no common suffix but
+ # still exactly one match, then select the
+ # first. (It could be that we have a completion
+ # which does * expansion, like '*.py', with
+ # exactly one match.)
+ if len(completions) == 1:
+ self.go_to_completion(0)
+
+ else:
+ # If the last operation was an insert, (not a delete), restart
+ # the completion coroutine.
+
+ if self.document.text_before_cursor == document.text_before_cursor:
+ return # Nothing changed.
+
+ if self.document.text_before_cursor.startswith(
+ document.text_before_cursor
+ ):
+ raise _Retry
+
+ return async_completer
+
+ def _create_auto_suggest_coroutine(self) -> Callable[[], Coroutine[Any, Any, None]]:
+ """
+ Create function for asynchronous auto suggestion.
+ (This can be in another thread.)
+ """
+
+ @_only_one_at_a_time
+ async def async_suggestor() -> None:
+ document = self.document
+
+ # Don't suggest when we already have a suggestion.
+ if self.suggestion or not self.auto_suggest:
+ return
+
+ suggestion = await self.auto_suggest.get_suggestion_async(self, document)
+
+ # Set suggestion only if the text was not yet changed.
+ if self.document == document:
+ # Set suggestion and redraw interface.
+ self.suggestion = suggestion
+ self.on_suggestion_set.fire()
+ else:
+ # Otherwise, restart thread.
+ raise _Retry
+
+ return async_suggestor
+
+ def _create_auto_validate_coroutine(
+ self,
+ ) -> Callable[[], Coroutine[Any, Any, None]]:
+ """
+ Create a function for asynchronous validation while typing.
+ (This can be in another thread.)
+ """
+
+ @_only_one_at_a_time
+ async def async_validator() -> None:
+ await self._validate_async()
+
+ return async_validator
+
+ def validate_and_handle(self) -> None:
+ """
+ Validate buffer and handle the accept action.
+ """
+ valid = self.validate(set_cursor=True)
+
+ # When the validation succeeded, accept the input.
+ if valid:
+ if self.accept_handler:
+ keep_text = self.accept_handler(self)
+ else:
+ keep_text = False
+
+ self.append_to_history()
+
+ if not keep_text:
+ self.reset()
+
+
+_T = TypeVar("_T", bound=Callable[..., Coroutine[Any, Any, None]])
+
+
+def _only_one_at_a_time(coroutine: _T) -> _T:
+ """
+ Decorator that only starts the coroutine only if the previous call has
+ finished. (Used to make sure that we have only one autocompleter, auto
+ suggestor and validator running at a time.)
+
+ When the coroutine raises `_Retry`, it is restarted.
+ """
+ running = False
+
+ @wraps(coroutine)
+ async def new_coroutine(*a: Any, **kw: Any) -> Any:
+ nonlocal running
+
+ # Don't start a new function, if the previous is still in progress.
+ if running:
+ return
+
+ running = True
+
+ try:
+ while True:
+ try:
+ await coroutine(*a, **kw)
+ except _Retry:
+ continue
+ else:
+ return None
+ finally:
+ running = False
+
+ return cast(_T, new_coroutine)
+
+
+class _Retry(Exception):
+ "Retry in `_only_one_at_a_time`."
+
+
+def indent(buffer: Buffer, from_row: int, to_row: int, count: int = 1) -> None:
+ """
+ Indent text of a :class:`.Buffer` object.
+ """
+ current_row = buffer.document.cursor_position_row
+ current_col = buffer.document.cursor_position_col
+ line_range = range(from_row, to_row)
+
+ # Apply transformation.
+ indent_content = " " * count
+ new_text = buffer.transform_lines(line_range, lambda l: indent_content + l)
+ buffer.document = Document(
+ new_text, Document(new_text).translate_row_col_to_index(current_row, 0)
+ )
+
+ # Place cursor in the same position in text after indenting
+ buffer.cursor_position += current_col + len(indent_content)
+
+
+def unindent(buffer: Buffer, from_row: int, to_row: int, count: int = 1) -> None:
+ """
+ Unindent text of a :class:`.Buffer` object.
+ """
+ current_row = buffer.document.cursor_position_row
+ current_col = buffer.document.cursor_position_col
+ line_range = range(from_row, to_row)
+
+ indent_content = " " * count
+
+ def transform(text: str) -> str:
+ remove = indent_content
+ if text.startswith(remove):
+ return text[len(remove) :]
+ else:
+ return text.lstrip()
+
+ # Apply transformation.
+ new_text = buffer.transform_lines(line_range, transform)
+ buffer.document = Document(
+ new_text, Document(new_text).translate_row_col_to_index(current_row, 0)
+ )
+
+ # Place cursor in the same position in text after dedent
+ buffer.cursor_position += current_col - len(indent_content)
+
+
+def reshape_text(buffer: Buffer, from_row: int, to_row: int) -> None:
+ """
+ Reformat text, taking the width into account.
+ `to_row` is included.
+ (Vi 'gq' operator.)
+ """
+ lines = buffer.text.splitlines(True)
+ lines_before = lines[:from_row]
+ lines_after = lines[to_row + 1 :]
+ lines_to_reformat = lines[from_row : to_row + 1]
+
+ if lines_to_reformat:
+ # Take indentation from the first line.
+ match = re.search(r"^\s*", lines_to_reformat[0])
+ length = match.end() if match else 0 # `match` can't be None, actually.
+
+ indent = lines_to_reformat[0][:length].replace("\n", "")
+
+ # Now, take all the 'words' from the lines to be reshaped.
+ words = "".join(lines_to_reformat).split()
+
+ # And reshape.
+ width = (buffer.text_width or 80) - len(indent)
+ reshaped_text = [indent]
+ current_width = 0
+ for w in words:
+ if current_width:
+ if len(w) + current_width + 1 > width:
+ reshaped_text.append("\n")
+ reshaped_text.append(indent)
+ current_width = 0
+ else:
+ reshaped_text.append(" ")
+ current_width += 1
+
+ reshaped_text.append(w)
+ current_width += len(w)
+
+ if reshaped_text[-1] != "\n":
+ reshaped_text.append("\n")
+
+ # Apply result.
+ buffer.document = Document(
+ text="".join(lines_before + reshaped_text + lines_after),
+ cursor_position=len("".join(lines_before + reshaped_text)),
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/cache.py b/lib/python3.12/site-packages/prompt_toolkit/cache.py
new file mode 100644
index 0000000000000000000000000000000000000000..01dd1f79d6582442b0932a6a252585a5d85aa0a5
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/cache.py
@@ -0,0 +1,127 @@
+from __future__ import annotations
+
+from collections import deque
+from functools import wraps
+from typing import Any, Callable, Dict, Generic, Hashable, Tuple, TypeVar, cast
+
+__all__ = [
+ "SimpleCache",
+ "FastDictCache",
+ "memoized",
+]
+
+_T = TypeVar("_T", bound=Hashable)
+_U = TypeVar("_U")
+
+
+class SimpleCache(Generic[_T, _U]):
+ """
+ Very simple cache that discards the oldest item when the cache size is
+ exceeded.
+
+ :param maxsize: Maximum size of the cache. (Don't make it too big.)
+ """
+
+ def __init__(self, maxsize: int = 8) -> None:
+ assert maxsize > 0
+
+ self._data: dict[_T, _U] = {}
+ self._keys: deque[_T] = deque()
+ self.maxsize: int = maxsize
+
+ def get(self, key: _T, getter_func: Callable[[], _U]) -> _U:
+ """
+ Get object from the cache.
+ If not found, call `getter_func` to resolve it, and put that on the top
+ of the cache instead.
+ """
+ # Look in cache first.
+ try:
+ return self._data[key]
+ except KeyError:
+ # Not found? Get it.
+ value = getter_func()
+ self._data[key] = value
+ self._keys.append(key)
+
+ # Remove the oldest key when the size is exceeded.
+ if len(self._data) > self.maxsize:
+ key_to_remove = self._keys.popleft()
+ if key_to_remove in self._data:
+ del self._data[key_to_remove]
+
+ return value
+
+ def clear(self) -> None:
+ "Clear cache."
+ self._data = {}
+ self._keys = deque()
+
+
+_K = TypeVar("_K", bound=Tuple[Hashable, ...])
+_V = TypeVar("_V")
+
+
+class FastDictCache(Dict[_K, _V]):
+ """
+ Fast, lightweight cache which keeps at most `size` items.
+ It will discard the oldest items in the cache first.
+
+ The cache is a dictionary, which doesn't keep track of access counts.
+ It is perfect to cache little immutable objects which are not expensive to
+ create, but where a dictionary lookup is still much faster than an object
+ instantiation.
+
+ :param get_value: Callable that's called in case of a missing key.
+ """
+
+ # NOTE: This cache is used to cache `prompt_toolkit.layout.screen.Char` and
+ # `prompt_toolkit.Document`. Make sure to keep this really lightweight.
+ # Accessing the cache should stay faster than instantiating new
+ # objects.
+ # (Dictionary lookups are really fast.)
+ # SimpleCache is still required for cases where the cache key is not
+ # the same as the arguments given to the function that creates the
+ # value.)
+ def __init__(self, get_value: Callable[..., _V], size: int = 1000000) -> None:
+ assert size > 0
+
+ self._keys: deque[_K] = deque()
+ self.get_value = get_value
+ self.size = size
+
+ def __missing__(self, key: _K) -> _V:
+ # Remove the oldest key when the size is exceeded.
+ if len(self) > self.size:
+ key_to_remove = self._keys.popleft()
+ if key_to_remove in self:
+ del self[key_to_remove]
+
+ result = self.get_value(*key)
+ self[key] = result
+ self._keys.append(key)
+ return result
+
+
+_F = TypeVar("_F", bound=Callable[..., object])
+
+
+def memoized(maxsize: int = 1024) -> Callable[[_F], _F]:
+ """
+ Memoization decorator for immutable classes and pure functions.
+ """
+
+ def decorator(obj: _F) -> _F:
+ cache: SimpleCache[Hashable, Any] = SimpleCache(maxsize=maxsize)
+
+ @wraps(obj)
+ def new_callable(*a: Any, **kw: Any) -> Any:
+ def create_new() -> Any:
+ return obj(*a, **kw)
+
+ key = (a, tuple(sorted(kw.items())))
+ return cache.get(key, create_new)
+
+ return cast(_F, new_callable)
+
+ return decorator
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..a8036e36bf243a0ab1824b9633c207b5816d81b5
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..172fe6f56774ae85437722ce1002f9b92e7089ea
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__init__.py
@@ -0,0 +1,5 @@
+from __future__ import annotations
+
+from .system import SystemCompleter
+
+__all__ = ["SystemCompleter"]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..46b7ae10feda84efe38b0fe9e02aa16321ba2f02
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__pycache__/system.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__pycache__/system.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..52230a868968daa4e8ece2bb0a40272ad52cf307
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/__pycache__/system.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/system.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/system.py
new file mode 100644
index 0000000000000000000000000000000000000000..5d990e5280567f313f211ac6c8de14622b755628
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/completers/system.py
@@ -0,0 +1,64 @@
+from __future__ import annotations
+
+from prompt_toolkit.completion.filesystem import ExecutableCompleter, PathCompleter
+from prompt_toolkit.contrib.regular_languages.compiler import compile
+from prompt_toolkit.contrib.regular_languages.completion import GrammarCompleter
+
+__all__ = [
+ "SystemCompleter",
+]
+
+
+class SystemCompleter(GrammarCompleter):
+ """
+ Completer for system commands.
+ """
+
+ def __init__(self) -> None:
+ # Compile grammar.
+ g = compile(
+ r"""
+ # First we have an executable.
+ (?P[^\s]+)
+
+ # Ignore literals in between.
+ (
+ \s+
+ ("[^"]*" | '[^']*' | [^'"]+ )
+ )*
+
+ \s+
+
+ # Filename as parameters.
+ (
+ (?P[^\s]+) |
+ "(?P[^\s]+)" |
+ '(?P[^\s]+)'
+ )
+ """,
+ escape_funcs={
+ "double_quoted_filename": (lambda string: string.replace('"', '\\"')),
+ "single_quoted_filename": (lambda string: string.replace("'", "\\'")),
+ },
+ unescape_funcs={
+ "double_quoted_filename": (
+ lambda string: string.replace('\\"', '"')
+ ), # XXX: not entirely correct.
+ "single_quoted_filename": (lambda string: string.replace("\\'", "'")),
+ },
+ )
+
+ # Create GrammarCompleter
+ super().__init__(
+ g,
+ {
+ "executable": ExecutableCompleter(),
+ "filename": PathCompleter(only_directories=False, expanduser=True),
+ "double_quoted_filename": PathCompleter(
+ only_directories=False, expanduser=True
+ ),
+ "single_quoted_filename": PathCompleter(
+ only_directories=False, expanduser=True
+ ),
+ },
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..38b027cd16f7a230b56a3375d70351ace31b6b9b
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__init__.py
@@ -0,0 +1,80 @@
+r"""
+Tool for expressing the grammar of an input as a regular language.
+==================================================================
+
+The grammar for the input of many simple command line interfaces can be
+expressed by a regular language. Examples are PDB (the Python debugger); a
+simple (bash-like) shell with "pwd", "cd", "cat" and "ls" commands; arguments
+that you can pass to an executable; etc. It is possible to use regular
+expressions for validation and parsing of such a grammar. (More about regular
+languages: http://en.wikipedia.org/wiki/Regular_language)
+
+Example
+-------
+
+Let's take the pwd/cd/cat/ls example. We want to have a shell that accepts
+these three commands. "cd" is followed by a quoted directory name and "cat" is
+followed by a quoted file name. (We allow quotes inside the filename when
+they're escaped with a backslash.) We could define the grammar using the
+following regular expression::
+
+ grammar = \s* (
+ pwd |
+ ls |
+ (cd \s+ " ([^"]|\.)+ ") |
+ (cat \s+ " ([^"]|\.)+ ")
+ ) \s*
+
+
+What can we do with this grammar?
+---------------------------------
+
+- Syntax highlighting: We could use this for instance to give file names
+ different color.
+- Parse the result: .. We can extract the file names and commands by using a
+ regular expression with named groups.
+- Input validation: .. Don't accept anything that does not match this grammar.
+ When combined with a parser, we can also recursively do
+ filename validation (and accept only existing files.)
+- Autocompletion: .... Each part of the grammar can have its own autocompleter.
+ "cat" has to be completed using file names, while "cd"
+ has to be completed using directory names.
+
+How does it work?
+-----------------
+
+As a user of this library, you have to define the grammar of the input as a
+regular expression. The parts of this grammar where autocompletion, validation
+or any other processing is required need to be marked using a regex named
+group. Like ``(?P...)`` for instance.
+
+When the input is processed for validation (for instance), the regex will
+execute, the named group is captured, and the validator associated with this
+named group will test the captured string.
+
+There is one tricky bit:
+
+ Often we operate on incomplete input (this is by definition the case for
+ autocompletion) and we have to decide for the cursor position in which
+ possible state the grammar it could be and in which way variables could be
+ matched up to that point.
+
+To solve this problem, the compiler takes the original regular expression and
+translates it into a set of other regular expressions which each match certain
+prefixes of the original regular expression. We generate one prefix regular
+expression for every named variable (with this variable being the end of that
+expression).
+
+
+TODO: some examples of:
+ - How to create a highlighter from this grammar.
+ - How to create a validator from this grammar.
+ - How to create an autocompleter from this grammar.
+ - How to create a parser from this grammar.
+"""
+
+from __future__ import annotations
+
+from .compiler import compile
+
+__all__ = ["compile"]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..32f9b869ff2167b22a18ff9924215d945ea6b7e1
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/compiler.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/compiler.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..97340987cf675b0b70b7fb9c5f672cbb77c61891
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/compiler.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/completion.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/completion.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..c2db8b82b5f7376b4269559265d09bbb81153363
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/completion.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/lexer.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/lexer.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..bedaff3d5f0619416f27aba93a28ce017b23d21a
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/lexer.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/regex_parser.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/regex_parser.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..4e9992fba04d6c48c7cb5e735276784acc18a1f7
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/regex_parser.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/validation.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/validation.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..1b0a4af0c75a7a38709218c23525e7eabde53fab
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/__pycache__/validation.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/compiler.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/compiler.py
new file mode 100644
index 0000000000000000000000000000000000000000..4009d54f2dc27e36bdd4d7469c0b368f19257c7b
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/compiler.py
@@ -0,0 +1,579 @@
+r"""
+Compiler for a regular grammar.
+
+Example usage::
+
+ # Create and compile grammar.
+ p = compile('add \s+ (?P[^\s]+) \s+ (?P[^\s]+)')
+
+ # Match input string.
+ m = p.match('add 23 432')
+
+ # Get variables.
+ m.variables().get('var1') # Returns "23"
+ m.variables().get('var2') # Returns "432"
+
+
+Partial matches are possible::
+
+ # Create and compile grammar.
+ p = compile('''
+ # Operators with two arguments.
+ ((?P[^\s]+) \s+ (?P[^\s]+) \s+ (?P[^\s]+)) |
+
+ # Operators with only one arguments.
+ ((?P[^\s]+) \s+ (?P[^\s]+))
+ ''')
+
+ # Match partial input string.
+ m = p.match_prefix('add 23')
+
+ # Get variables. (Notice that both operator1 and operator2 contain the
+ # value "add".) This is because our input is incomplete, and we don't know
+ # yet in which rule of the regex we we'll end up. It could also be that
+ # `operator1` and `operator2` have a different autocompleter and we want to
+ # call all possible autocompleters that would result in valid input.)
+ m.variables().get('var1') # Returns "23"
+ m.variables().get('operator1') # Returns "add"
+ m.variables().get('operator2') # Returns "add"
+
+"""
+
+from __future__ import annotations
+
+import re
+from typing import Callable, Dict, Iterable, Iterator, Pattern, TypeVar, overload
+from typing import Match as RegexMatch
+
+from .regex_parser import (
+ AnyNode,
+ Lookahead,
+ Node,
+ NodeSequence,
+ Regex,
+ Repeat,
+ Variable,
+ parse_regex,
+ tokenize_regex,
+)
+
+__all__ = ["compile", "Match", "Variables"]
+
+
+# Name of the named group in the regex, matching trailing input.
+# (Trailing input is when the input contains characters after the end of the
+# expression has been matched.)
+_INVALID_TRAILING_INPUT = "invalid_trailing"
+
+EscapeFuncDict = Dict[str, Callable[[str], str]]
+
+
+class _CompiledGrammar:
+ """
+ Compiles a grammar. This will take the parse tree of a regular expression
+ and compile the grammar.
+
+ :param root_node: :class~`.regex_parser.Node` instance.
+ :param escape_funcs: `dict` mapping variable names to escape callables.
+ :param unescape_funcs: `dict` mapping variable names to unescape callables.
+ """
+
+ def __init__(
+ self,
+ root_node: Node,
+ escape_funcs: EscapeFuncDict | None = None,
+ unescape_funcs: EscapeFuncDict | None = None,
+ ) -> None:
+ self.root_node = root_node
+ self.escape_funcs = escape_funcs or {}
+ self.unescape_funcs = unescape_funcs or {}
+
+ #: Dictionary that will map the regex names to Node instances.
+ self._group_names_to_nodes: dict[
+ str, str
+ ] = {} # Maps regex group names to varnames.
+ counter = [0]
+
+ def create_group_func(node: Variable) -> str:
+ name = f"n{counter[0]}"
+ self._group_names_to_nodes[name] = node.varname
+ counter[0] += 1
+ return name
+
+ # Compile regex strings.
+ self._re_pattern = f"^{self._transform(root_node, create_group_func)}$"
+ self._re_prefix_patterns = list(
+ self._transform_prefix(root_node, create_group_func)
+ )
+
+ # Compile the regex itself.
+ flags = re.DOTALL # Note that we don't need re.MULTILINE! (^ and $
+ # still represent the start and end of input text.)
+ self._re = re.compile(self._re_pattern, flags)
+ self._re_prefix = [re.compile(t, flags) for t in self._re_prefix_patterns]
+
+ # We compile one more set of regexes, similar to `_re_prefix`, but accept any trailing
+ # input. This will ensure that we can still highlight the input correctly, even when the
+ # input contains some additional characters at the end that don't match the grammar.)
+ self._re_prefix_with_trailing_input = [
+ re.compile(
+ r"(?:{})(?P<{}>.*?)$".format(t.rstrip("$"), _INVALID_TRAILING_INPUT),
+ flags,
+ )
+ for t in self._re_prefix_patterns
+ ]
+
+ def escape(self, varname: str, value: str) -> str:
+ """
+ Escape `value` to fit in the place of this variable into the grammar.
+ """
+ f = self.escape_funcs.get(varname)
+ return f(value) if f else value
+
+ def unescape(self, varname: str, value: str) -> str:
+ """
+ Unescape `value`.
+ """
+ f = self.unescape_funcs.get(varname)
+ return f(value) if f else value
+
+ @classmethod
+ def _transform(
+ cls, root_node: Node, create_group_func: Callable[[Variable], str]
+ ) -> str:
+ """
+ Turn a :class:`Node` object into a regular expression.
+
+ :param root_node: The :class:`Node` instance for which we generate the grammar.
+ :param create_group_func: A callable which takes a `Node` and returns the next
+ free name for this node.
+ """
+
+ def transform(node: Node) -> str:
+ # Turn `AnyNode` into an OR.
+ if isinstance(node, AnyNode):
+ return "(?:{})".format("|".join(transform(c) for c in node.children))
+
+ # Concatenate a `NodeSequence`
+ elif isinstance(node, NodeSequence):
+ return "".join(transform(c) for c in node.children)
+
+ # For Regex and Lookahead nodes, just insert them literally.
+ elif isinstance(node, Regex):
+ return node.regex
+
+ elif isinstance(node, Lookahead):
+ before = "(?!" if node.negative else "(="
+ return before + transform(node.childnode) + ")"
+
+ # A `Variable` wraps the children into a named group.
+ elif isinstance(node, Variable):
+ return f"(?P<{create_group_func(node)}>{transform(node.childnode)})"
+
+ # `Repeat`.
+ elif isinstance(node, Repeat):
+ if node.max_repeat is None:
+ if node.min_repeat == 0:
+ repeat_sign = "*"
+ elif node.min_repeat == 1:
+ repeat_sign = "+"
+ else:
+ repeat_sign = "{%i,%s}" % (
+ node.min_repeat,
+ ("" if node.max_repeat is None else str(node.max_repeat)),
+ )
+
+ return "(?:{}){}{}".format(
+ transform(node.childnode),
+ repeat_sign,
+ ("" if node.greedy else "?"),
+ )
+ else:
+ raise TypeError(f"Got {node!r}")
+
+ return transform(root_node)
+
+ @classmethod
+ def _transform_prefix(
+ cls, root_node: Node, create_group_func: Callable[[Variable], str]
+ ) -> Iterable[str]:
+ """
+ Yield all the regular expressions matching a prefix of the grammar
+ defined by the `Node` instance.
+
+ For each `Variable`, one regex pattern will be generated, with this
+ named group at the end. This is required because a regex engine will
+ terminate once a match is found. For autocompletion however, we need
+ the matches for all possible paths, so that we can provide completions
+ for each `Variable`.
+
+ - So, in the case of an `Any` (`A|B|C)', we generate a pattern for each
+ clause. This is one for `A`, one for `B` and one for `C`. Unless some
+ groups don't contain a `Variable`, then these can be merged together.
+ - In the case of a `NodeSequence` (`ABC`), we generate a pattern for
+ each prefix that ends with a variable, and one pattern for the whole
+ sequence. So, that's one for `A`, one for `AB` and one for `ABC`.
+
+ :param root_node: The :class:`Node` instance for which we generate the grammar.
+ :param create_group_func: A callable which takes a `Node` and returns the next
+ free name for this node.
+ """
+
+ def contains_variable(node: Node) -> bool:
+ if isinstance(node, Regex):
+ return False
+ elif isinstance(node, Variable):
+ return True
+ elif isinstance(node, (Lookahead, Repeat)):
+ return contains_variable(node.childnode)
+ elif isinstance(node, (NodeSequence, AnyNode)):
+ return any(contains_variable(child) for child in node.children)
+
+ return False
+
+ def transform(node: Node) -> Iterable[str]:
+ # Generate separate pattern for all terms that contain variables
+ # within this OR. Terms that don't contain a variable can be merged
+ # together in one pattern.
+ if isinstance(node, AnyNode):
+ # If we have a definition like:
+ # (?P .*) | (?P .*)
+ # Then we want to be able to generate completions for both the
+ # name as well as the city. We do this by yielding two
+ # different regular expressions, because the engine won't
+ # follow multiple paths, if multiple are possible.
+ children_with_variable = []
+ children_without_variable = []
+ for c in node.children:
+ if contains_variable(c):
+ children_with_variable.append(c)
+ else:
+ children_without_variable.append(c)
+
+ for c in children_with_variable:
+ yield from transform(c)
+
+ # Merge options without variable together.
+ if children_without_variable:
+ yield "|".join(
+ r for c in children_without_variable for r in transform(c)
+ )
+
+ # For a sequence, generate a pattern for each prefix that ends with
+ # a variable + one pattern of the complete sequence.
+ # (This is because, for autocompletion, we match the text before
+ # the cursor, and completions are given for the variable that we
+ # match right before the cursor.)
+ elif isinstance(node, NodeSequence):
+ # For all components in the sequence, compute prefix patterns,
+ # as well as full patterns.
+ complete = [cls._transform(c, create_group_func) for c in node.children]
+ prefixes = [list(transform(c)) for c in node.children]
+ variable_nodes = [contains_variable(c) for c in node.children]
+
+ # If any child is contains a variable, we should yield a
+ # pattern up to that point, so that we are sure this will be
+ # matched.
+ for i in range(len(node.children)):
+ if variable_nodes[i]:
+ for c_str in prefixes[i]:
+ yield "".join(complete[:i]) + c_str
+
+ # If there are non-variable nodes, merge all the prefixes into
+ # one pattern. If the input is: "[part1] [part2] [part3]", then
+ # this gets compiled into:
+ # (complete1 + (complete2 + (complete3 | partial3) | partial2) | partial1 )
+ # For nodes that contain a variable, we skip the "|partial"
+ # part here, because thees are matched with the previous
+ # patterns.
+ if not all(variable_nodes):
+ result = []
+
+ # Start with complete patterns.
+ for i in range(len(node.children)):
+ result.append("(?:")
+ result.append(complete[i])
+
+ # Add prefix patterns.
+ for i in range(len(node.children) - 1, -1, -1):
+ if variable_nodes[i]:
+ # No need to yield a prefix for this one, we did
+ # the variable prefixes earlier.
+ result.append(")")
+ else:
+ result.append("|(?:")
+ # If this yields multiple, we should yield all combinations.
+ assert len(prefixes[i]) == 1
+ result.append(prefixes[i][0])
+ result.append("))")
+
+ yield "".join(result)
+
+ elif isinstance(node, Regex):
+ yield f"(?:{node.regex})?"
+
+ elif isinstance(node, Lookahead):
+ if node.negative:
+ yield f"(?!{cls._transform(node.childnode, create_group_func)})"
+ else:
+ # Not sure what the correct semantics are in this case.
+ # (Probably it's not worth implementing this.)
+ raise Exception("Positive lookahead not yet supported.")
+
+ elif isinstance(node, Variable):
+ # (Note that we should not append a '?' here. the 'transform'
+ # method will already recursively do that.)
+ for c_str in transform(node.childnode):
+ yield f"(?P<{create_group_func(node)}>{c_str})"
+
+ elif isinstance(node, Repeat):
+ # If we have a repetition of 8 times. That would mean that the
+ # current input could have for instance 7 times a complete
+ # match, followed by a partial match.
+ prefix = cls._transform(node.childnode, create_group_func)
+
+ if node.max_repeat == 1:
+ yield from transform(node.childnode)
+ else:
+ for c_str in transform(node.childnode):
+ if node.max_repeat:
+ repeat_sign = "{,%i}" % (node.max_repeat - 1)
+ else:
+ repeat_sign = "*"
+ yield "(?:{}){}{}{}".format(
+ prefix,
+ repeat_sign,
+ ("" if node.greedy else "?"),
+ c_str,
+ )
+
+ else:
+ raise TypeError(f"Got {node!r}")
+
+ for r in transform(root_node):
+ yield f"^(?:{r})$"
+
+ def match(self, string: str) -> Match | None:
+ """
+ Match the string with the grammar.
+ Returns a :class:`Match` instance or `None` when the input doesn't match the grammar.
+
+ :param string: The input string.
+ """
+ m = self._re.match(string)
+
+ if m:
+ return Match(
+ string, [(self._re, m)], self._group_names_to_nodes, self.unescape_funcs
+ )
+ return None
+
+ def match_prefix(self, string: str) -> Match | None:
+ """
+ Do a partial match of the string with the grammar. The returned
+ :class:`Match` instance can contain multiple representations of the
+ match. This will never return `None`. If it doesn't match at all, the "trailing input"
+ part will capture all of the input.
+
+ :param string: The input string.
+ """
+ # First try to match using `_re_prefix`. If nothing is found, use the patterns that
+ # also accept trailing characters.
+ for patterns in [self._re_prefix, self._re_prefix_with_trailing_input]:
+ matches = [(r, r.match(string)) for r in patterns]
+ matches2 = [(r, m) for r, m in matches if m]
+
+ if matches2 != []:
+ return Match(
+ string, matches2, self._group_names_to_nodes, self.unescape_funcs
+ )
+
+ return None
+
+
+class Match:
+ """
+ :param string: The input string.
+ :param re_matches: List of (compiled_re_pattern, re_match) tuples.
+ :param group_names_to_nodes: Dictionary mapping all the re group names to the matching Node instances.
+ """
+
+ def __init__(
+ self,
+ string: str,
+ re_matches: list[tuple[Pattern[str], RegexMatch[str]]],
+ group_names_to_nodes: dict[str, str],
+ unescape_funcs: dict[str, Callable[[str], str]],
+ ):
+ self.string = string
+ self._re_matches = re_matches
+ self._group_names_to_nodes = group_names_to_nodes
+ self._unescape_funcs = unescape_funcs
+
+ def _nodes_to_regs(self) -> list[tuple[str, tuple[int, int]]]:
+ """
+ Return a list of (varname, reg) tuples.
+ """
+
+ def get_tuples() -> Iterable[tuple[str, tuple[int, int]]]:
+ for r, re_match in self._re_matches:
+ for group_name, group_index in r.groupindex.items():
+ if group_name != _INVALID_TRAILING_INPUT:
+ regs = re_match.regs
+ reg = regs[group_index]
+ node = self._group_names_to_nodes[group_name]
+ yield (node, reg)
+
+ return list(get_tuples())
+
+ def _nodes_to_values(self) -> list[tuple[str, str, tuple[int, int]]]:
+ """
+ Returns list of (Node, string_value) tuples.
+ """
+
+ def is_none(sl: tuple[int, int]) -> bool:
+ return sl[0] == -1 and sl[1] == -1
+
+ def get(sl: tuple[int, int]) -> str:
+ return self.string[sl[0] : sl[1]]
+
+ return [
+ (varname, get(slice), slice)
+ for varname, slice in self._nodes_to_regs()
+ if not is_none(slice)
+ ]
+
+ def _unescape(self, varname: str, value: str) -> str:
+ unwrapper = self._unescape_funcs.get(varname)
+ return unwrapper(value) if unwrapper else value
+
+ def variables(self) -> Variables:
+ """
+ Returns :class:`Variables` instance.
+ """
+ return Variables(
+ [(k, self._unescape(k, v), sl) for k, v, sl in self._nodes_to_values()]
+ )
+
+ def trailing_input(self) -> MatchVariable | None:
+ """
+ Get the `MatchVariable` instance, representing trailing input, if there is any.
+ "Trailing input" is input at the end that does not match the grammar anymore, but
+ when this is removed from the end of the input, the input would be a valid string.
+ """
+ slices: list[tuple[int, int]] = []
+
+ # Find all regex group for the name _INVALID_TRAILING_INPUT.
+ for r, re_match in self._re_matches:
+ for group_name, group_index in r.groupindex.items():
+ if group_name == _INVALID_TRAILING_INPUT:
+ slices.append(re_match.regs[group_index])
+
+ # Take the smallest part. (Smaller trailing text means that a larger input has
+ # been matched, so that is better.)
+ if slices:
+ slice = (max(i[0] for i in slices), max(i[1] for i in slices))
+ value = self.string[slice[0] : slice[1]]
+ return MatchVariable("", value, slice)
+ return None
+
+ def end_nodes(self) -> Iterable[MatchVariable]:
+ """
+ Yields `MatchVariable` instances for all the nodes having their end
+ position at the end of the input string.
+ """
+ for varname, reg in self._nodes_to_regs():
+ # If this part goes until the end of the input string.
+ if reg[1] == len(self.string):
+ value = self._unescape(varname, self.string[reg[0] : reg[1]])
+ yield MatchVariable(varname, value, (reg[0], reg[1]))
+
+
+_T = TypeVar("_T")
+
+
+class Variables:
+ def __init__(self, tuples: list[tuple[str, str, tuple[int, int]]]) -> None:
+ #: List of (varname, value, slice) tuples.
+ self._tuples = tuples
+
+ def __repr__(self) -> str:
+ return "{}({})".format(
+ self.__class__.__name__,
+ ", ".join(f"{k}={v!r}" for k, v, _ in self._tuples),
+ )
+
+ @overload
+ def get(self, key: str) -> str | None: ...
+
+ @overload
+ def get(self, key: str, default: str | _T) -> str | _T: ...
+
+ def get(self, key: str, default: str | _T | None = None) -> str | _T | None:
+ items = self.getall(key)
+ return items[0] if items else default
+
+ def getall(self, key: str) -> list[str]:
+ return [v for k, v, _ in self._tuples if k == key]
+
+ def __getitem__(self, key: str) -> str | None:
+ return self.get(key)
+
+ def __iter__(self) -> Iterator[MatchVariable]:
+ """
+ Yield `MatchVariable` instances.
+ """
+ for varname, value, slice in self._tuples:
+ yield MatchVariable(varname, value, slice)
+
+
+class MatchVariable:
+ """
+ Represents a match of a variable in the grammar.
+
+ :param varname: (string) Name of the variable.
+ :param value: (string) Value of this variable.
+ :param slice: (start, stop) tuple, indicating the position of this variable
+ in the input string.
+ """
+
+ def __init__(self, varname: str, value: str, slice: tuple[int, int]) -> None:
+ self.varname = varname
+ self.value = value
+ self.slice = slice
+
+ self.start = self.slice[0]
+ self.stop = self.slice[1]
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.varname!r}, {self.value!r})"
+
+
+def compile(
+ expression: str,
+ escape_funcs: EscapeFuncDict | None = None,
+ unescape_funcs: EscapeFuncDict | None = None,
+) -> _CompiledGrammar:
+ """
+ Compile grammar (given as regex string), returning a `CompiledGrammar`
+ instance.
+ """
+ return _compile_from_parse_tree(
+ parse_regex(tokenize_regex(expression)),
+ escape_funcs=escape_funcs,
+ unescape_funcs=unescape_funcs,
+ )
+
+
+def _compile_from_parse_tree(
+ root_node: Node,
+ escape_funcs: EscapeFuncDict | None = None,
+ unescape_funcs: EscapeFuncDict | None = None,
+) -> _CompiledGrammar:
+ """
+ Compile grammar (given as parse tree), returning a `CompiledGrammar`
+ instance.
+ """
+ return _CompiledGrammar(
+ root_node, escape_funcs=escape_funcs, unescape_funcs=unescape_funcs
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/completion.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/completion.py
new file mode 100644
index 0000000000000000000000000000000000000000..19ebaad9389064dec387d0b325f4a310aee21c8f
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/completion.py
@@ -0,0 +1,100 @@
+"""
+Completer for a regular grammar.
+"""
+
+from __future__ import annotations
+
+from typing import Iterable
+
+from prompt_toolkit.completion import CompleteEvent, Completer, Completion
+from prompt_toolkit.document import Document
+
+from .compiler import Match, _CompiledGrammar
+
+__all__ = [
+ "GrammarCompleter",
+]
+
+
+class GrammarCompleter(Completer):
+ """
+ Completer which can be used for autocompletion according to variables in
+ the grammar. Each variable can have a different autocompleter.
+
+ :param compiled_grammar: `GrammarCompleter` instance.
+ :param completers: `dict` mapping variable names of the grammar to the
+ `Completer` instances to be used for each variable.
+ """
+
+ def __init__(
+ self, compiled_grammar: _CompiledGrammar, completers: dict[str, Completer]
+ ) -> None:
+ self.compiled_grammar = compiled_grammar
+ self.completers = completers
+
+ def get_completions(
+ self, document: Document, complete_event: CompleteEvent
+ ) -> Iterable[Completion]:
+ m = self.compiled_grammar.match_prefix(document.text_before_cursor)
+
+ if m:
+ yield from self._remove_duplicates(
+ self._get_completions_for_match(m, complete_event)
+ )
+
+ def _get_completions_for_match(
+ self, match: Match, complete_event: CompleteEvent
+ ) -> Iterable[Completion]:
+ """
+ Yield all the possible completions for this input string.
+ (The completer assumes that the cursor position was at the end of the
+ input string.)
+ """
+ for match_variable in match.end_nodes():
+ varname = match_variable.varname
+ start = match_variable.start
+
+ completer = self.completers.get(varname)
+
+ if completer:
+ text = match_variable.value
+
+ # Unwrap text.
+ unwrapped_text = self.compiled_grammar.unescape(varname, text)
+
+ # Create a document, for the completions API (text/cursor_position)
+ document = Document(unwrapped_text, len(unwrapped_text))
+
+ # Call completer
+ for completion in completer.get_completions(document, complete_event):
+ new_text = (
+ unwrapped_text[: len(text) + completion.start_position]
+ + completion.text
+ )
+
+ # Wrap again.
+ yield Completion(
+ text=self.compiled_grammar.escape(varname, new_text),
+ start_position=start - len(match.string),
+ display=completion.display,
+ display_meta=completion.display_meta,
+ )
+
+ def _remove_duplicates(self, items: Iterable[Completion]) -> Iterable[Completion]:
+ """
+ Remove duplicates, while keeping the order.
+ (Sometimes we have duplicates, because the there several matches of the
+ same grammar, each yielding similar completions.)
+ """
+
+ def hash_completion(completion: Completion) -> tuple[str, int]:
+ return completion.text, completion.start_position
+
+ yielded_so_far: set[tuple[str, int]] = set()
+
+ for completion in items:
+ hash_value = hash_completion(completion)
+
+ if hash_value not in yielded_so_far:
+ yielded_so_far.add(hash_value)
+ yield completion
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/lexer.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/lexer.py
new file mode 100644
index 0000000000000000000000000000000000000000..c5434cfdee07c90cf082a6117bc66a08a2b17f83
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/lexer.py
@@ -0,0 +1,94 @@
+"""
+`GrammarLexer` is compatible with other lexers and can be used to highlight
+the input using a regular grammar with annotations.
+"""
+
+from __future__ import annotations
+
+from typing import Callable
+
+from prompt_toolkit.document import Document
+from prompt_toolkit.formatted_text.base import StyleAndTextTuples
+from prompt_toolkit.formatted_text.utils import split_lines
+from prompt_toolkit.lexers import Lexer
+
+from .compiler import _CompiledGrammar
+
+__all__ = [
+ "GrammarLexer",
+]
+
+
+class GrammarLexer(Lexer):
+ """
+ Lexer which can be used for highlighting of fragments according to variables in the grammar.
+
+ (It does not actual lexing of the string, but it exposes an API, compatible
+ with the Pygments lexer class.)
+
+ :param compiled_grammar: Grammar as returned by the `compile()` function.
+ :param lexers: Dictionary mapping variable names of the regular grammar to
+ the lexers that should be used for this part. (This can
+ call other lexers recursively.) If you wish a part of the
+ grammar to just get one fragment, use a
+ `prompt_toolkit.lexers.SimpleLexer`.
+ """
+
+ def __init__(
+ self,
+ compiled_grammar: _CompiledGrammar,
+ default_style: str = "",
+ lexers: dict[str, Lexer] | None = None,
+ ) -> None:
+ self.compiled_grammar = compiled_grammar
+ self.default_style = default_style
+ self.lexers = lexers or {}
+
+ def _get_text_fragments(self, text: str) -> StyleAndTextTuples:
+ m = self.compiled_grammar.match_prefix(text)
+
+ if m:
+ characters: StyleAndTextTuples = [(self.default_style, c) for c in text]
+
+ for v in m.variables():
+ # If we have a `Lexer` instance for this part of the input.
+ # Tokenize recursively and apply tokens.
+ lexer = self.lexers.get(v.varname)
+
+ if lexer:
+ document = Document(text[v.start : v.stop])
+ lexer_tokens_for_line = lexer.lex_document(document)
+ text_fragments: StyleAndTextTuples = []
+ for i in range(len(document.lines)):
+ text_fragments.extend(lexer_tokens_for_line(i))
+ text_fragments.append(("", "\n"))
+ if text_fragments:
+ text_fragments.pop()
+
+ i = v.start
+ for t, s, *_ in text_fragments:
+ for c in s:
+ if characters[i][0] == self.default_style:
+ characters[i] = (t, characters[i][1])
+ i += 1
+
+ # Highlight trailing input.
+ trailing_input = m.trailing_input()
+ if trailing_input:
+ for i in range(trailing_input.start, trailing_input.stop):
+ characters[i] = ("class:trailing-input", characters[i][1])
+
+ return characters
+ else:
+ return [("", text)]
+
+ def lex_document(self, document: Document) -> Callable[[int], StyleAndTextTuples]:
+ lines = list(split_lines(self._get_text_fragments(document.text)))
+
+ def get_line(lineno: int) -> StyleAndTextTuples:
+ try:
+ return lines[lineno]
+ except IndexError:
+ return []
+
+ return get_line
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/regex_parser.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/regex_parser.py
new file mode 100644
index 0000000000000000000000000000000000000000..353e54ff726e7a5d16bec7a0d50d828eb7880b1e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/regex_parser.py
@@ -0,0 +1,279 @@
+"""
+Parser for parsing a regular expression.
+Take a string representing a regular expression and return the root node of its
+parse tree.
+
+usage::
+
+ root_node = parse_regex('(hello|world)')
+
+Remarks:
+- The regex parser processes multiline, it ignores all whitespace and supports
+ multiple named groups with the same name and #-style comments.
+
+Limitations:
+- Lookahead is not supported.
+"""
+
+from __future__ import annotations
+
+import re
+
+__all__ = [
+ "Repeat",
+ "Variable",
+ "Regex",
+ "Lookahead",
+ "tokenize_regex",
+ "parse_regex",
+]
+
+
+class Node:
+ """
+ Base class for all the grammar nodes.
+ (You don't initialize this one.)
+ """
+
+ def __add__(self, other_node: Node) -> NodeSequence:
+ return NodeSequence([self, other_node])
+
+ def __or__(self, other_node: Node) -> AnyNode:
+ return AnyNode([self, other_node])
+
+
+class AnyNode(Node):
+ """
+ Union operation (OR operation) between several grammars. You don't
+ initialize this yourself, but it's a result of a "Grammar1 | Grammar2"
+ operation.
+ """
+
+ def __init__(self, children: list[Node]) -> None:
+ self.children = children
+
+ def __or__(self, other_node: Node) -> AnyNode:
+ return AnyNode(self.children + [other_node])
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.children!r})"
+
+
+class NodeSequence(Node):
+ """
+ Concatenation operation of several grammars. You don't initialize this
+ yourself, but it's a result of a "Grammar1 + Grammar2" operation.
+ """
+
+ def __init__(self, children: list[Node]) -> None:
+ self.children = children
+
+ def __add__(self, other_node: Node) -> NodeSequence:
+ return NodeSequence(self.children + [other_node])
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.children!r})"
+
+
+class Regex(Node):
+ """
+ Regular expression.
+ """
+
+ def __init__(self, regex: str) -> None:
+ re.compile(regex) # Validate
+
+ self.regex = regex
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(/{self.regex}/)"
+
+
+class Lookahead(Node):
+ """
+ Lookahead expression.
+ """
+
+ def __init__(self, childnode: Node, negative: bool = False) -> None:
+ self.childnode = childnode
+ self.negative = negative
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.childnode!r})"
+
+
+class Variable(Node):
+ """
+ Mark a variable in the regular grammar. This will be translated into a
+ named group. Each variable can have his own completer, validator, etc..
+
+ :param childnode: The grammar which is wrapped inside this variable.
+ :param varname: String.
+ """
+
+ def __init__(self, childnode: Node, varname: str = "") -> None:
+ self.childnode = childnode
+ self.varname = varname
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(childnode={self.childnode!r}, varname={self.varname!r})"
+
+
+class Repeat(Node):
+ def __init__(
+ self,
+ childnode: Node,
+ min_repeat: int = 0,
+ max_repeat: int | None = None,
+ greedy: bool = True,
+ ) -> None:
+ self.childnode = childnode
+ self.min_repeat = min_repeat
+ self.max_repeat = max_repeat
+ self.greedy = greedy
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(childnode={self.childnode!r})"
+
+
+def tokenize_regex(input: str) -> list[str]:
+ """
+ Takes a string, representing a regular expression as input, and tokenizes
+ it.
+
+ :param input: string, representing a regular expression.
+ :returns: List of tokens.
+ """
+ # Regular expression for tokenizing other regular expressions.
+ p = re.compile(
+ r"""^(
+ \(\?P\<[a-zA-Z0-9_-]+\> | # Start of named group.
+ \(\?#[^)]*\) | # Comment
+ \(\?= | # Start of lookahead assertion
+ \(\?! | # Start of negative lookahead assertion
+ \(\?<= | # If preceded by.
+ \(\?< | # If not preceded by.
+ \(?: | # Start of group. (non capturing.)
+ \( | # Start of group.
+ \(?[iLmsux] | # Flags.
+ \(?P=[a-zA-Z]+\) | # Back reference to named group
+ \) | # End of group.
+ \{[^{}]*\} | # Repetition
+ \*\? | \+\? | \?\?\ | # Non greedy repetition.
+ \* | \+ | \? | # Repetition
+ \#.*\n | # Comment
+ \\. |
+
+ # Character group.
+ \[
+ ( [^\]\\] | \\.)*
+ \] |
+
+ [^(){}] |
+ .
+ )""",
+ re.VERBOSE,
+ )
+
+ tokens = []
+
+ while input:
+ m = p.match(input)
+ if m:
+ token, input = input[: m.end()], input[m.end() :]
+ if not token.isspace():
+ tokens.append(token)
+ else:
+ raise Exception("Could not tokenize input regex.")
+
+ return tokens
+
+
+def parse_regex(regex_tokens: list[str]) -> Node:
+ """
+ Takes a list of tokens from the tokenizer, and returns a parse tree.
+ """
+ # We add a closing brace because that represents the final pop of the stack.
+ tokens: list[str] = [")"] + regex_tokens[::-1]
+
+ def wrap(lst: list[Node]) -> Node:
+ """Turn list into sequence when it contains several items."""
+ if len(lst) == 1:
+ return lst[0]
+ else:
+ return NodeSequence(lst)
+
+ def _parse() -> Node:
+ or_list: list[list[Node]] = []
+ result: list[Node] = []
+
+ def wrapped_result() -> Node:
+ if or_list == []:
+ return wrap(result)
+ else:
+ or_list.append(result)
+ return AnyNode([wrap(i) for i in or_list])
+
+ while tokens:
+ t = tokens.pop()
+
+ if t.startswith("(?P<"):
+ variable = Variable(_parse(), varname=t[4:-1])
+ result.append(variable)
+
+ elif t in ("*", "*?"):
+ greedy = t == "*"
+ result[-1] = Repeat(result[-1], greedy=greedy)
+
+ elif t in ("+", "+?"):
+ greedy = t == "+"
+ result[-1] = Repeat(result[-1], min_repeat=1, greedy=greedy)
+
+ elif t in ("?", "??"):
+ if result == []:
+ raise Exception("Nothing to repeat." + repr(tokens))
+ else:
+ greedy = t == "?"
+ result[-1] = Repeat(
+ result[-1], min_repeat=0, max_repeat=1, greedy=greedy
+ )
+
+ elif t == "|":
+ or_list.append(result)
+ result = []
+
+ elif t in ("(", "(?:"):
+ result.append(_parse())
+
+ elif t == "(?!":
+ result.append(Lookahead(_parse(), negative=True))
+
+ elif t == "(?=":
+ result.append(Lookahead(_parse(), negative=False))
+
+ elif t == ")":
+ return wrapped_result()
+
+ elif t.startswith("#"):
+ pass
+
+ elif t.startswith("{"):
+ # TODO: implement!
+ raise Exception(f"{t}-style repetition not yet supported")
+
+ elif t.startswith("(?"):
+ raise Exception(f"{t!r} not supported")
+
+ elif t.isspace():
+ pass
+ else:
+ result.append(Regex(t))
+
+ raise Exception("Expecting ')' token")
+
+ result = _parse()
+
+ if len(tokens) != 0:
+ raise Exception("Unmatched parentheses.")
+ else:
+ return result
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/validation.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/validation.py
new file mode 100644
index 0000000000000000000000000000000000000000..e6cfd741328f293bee89e88f0524773c1d724a7e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/regular_languages/validation.py
@@ -0,0 +1,60 @@
+"""
+Validator for a regular language.
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit.document import Document
+from prompt_toolkit.validation import ValidationError, Validator
+
+from .compiler import _CompiledGrammar
+
+__all__ = [
+ "GrammarValidator",
+]
+
+
+class GrammarValidator(Validator):
+ """
+ Validator which can be used for validation according to variables in
+ the grammar. Each variable can have its own validator.
+
+ :param compiled_grammar: `GrammarCompleter` instance.
+ :param validators: `dict` mapping variable names of the grammar to the
+ `Validator` instances to be used for each variable.
+ """
+
+ def __init__(
+ self, compiled_grammar: _CompiledGrammar, validators: dict[str, Validator]
+ ) -> None:
+ self.compiled_grammar = compiled_grammar
+ self.validators = validators
+
+ def validate(self, document: Document) -> None:
+ # Parse input document.
+ # We use `match`, not `match_prefix`, because for validation, we want
+ # the actual, unambiguous interpretation of the input.
+ m = self.compiled_grammar.match(document.text)
+
+ if m:
+ for v in m.variables():
+ validator = self.validators.get(v.varname)
+
+ if validator:
+ # Unescape text.
+ unwrapped_text = self.compiled_grammar.unescape(v.varname, v.value)
+
+ # Create a document, for the completions API (text/cursor_position)
+ inner_document = Document(unwrapped_text, len(unwrapped_text))
+
+ try:
+ validator.validate(inner_document)
+ except ValidationError as e:
+ raise ValidationError(
+ cursor_position=v.start + e.cursor_position,
+ message=e.message,
+ ) from e
+ else:
+ raise ValidationError(
+ cursor_position=len(document.text), message="Invalid command"
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..bbc1c215b1748a27518f98a2d25d03a45494baf4
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__init__.py
@@ -0,0 +1,8 @@
+from __future__ import annotations
+
+from .server import PromptToolkitSSHServer, PromptToolkitSSHSession
+
+__all__ = [
+ "PromptToolkitSSHSession",
+ "PromptToolkitSSHServer",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..e5fb2b0798685a2dc7ca1eea3b21f41a99b01e38
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__pycache__/server.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__pycache__/server.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..0379cf3f7e6b30a4ae1fd9305a486da61fae0eee
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/__pycache__/server.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/server.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/server.py
new file mode 100644
index 0000000000000000000000000000000000000000..4badc6cad04b6144b7ff32a1f88b54b9e4bbedd4
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/ssh/server.py
@@ -0,0 +1,178 @@
+"""
+Utility for running a prompt_toolkit application in an asyncssh server.
+"""
+
+from __future__ import annotations
+
+import asyncio
+import traceback
+from asyncio import get_running_loop
+from typing import Any, Callable, Coroutine, TextIO, cast
+
+import asyncssh
+
+from prompt_toolkit.application.current import AppSession, create_app_session
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.input import PipeInput, create_pipe_input
+from prompt_toolkit.output.vt100 import Vt100_Output
+
+__all__ = ["PromptToolkitSSHSession", "PromptToolkitSSHServer"]
+
+
+class PromptToolkitSSHSession(asyncssh.SSHServerSession): # type: ignore
+ def __init__(
+ self,
+ interact: Callable[[PromptToolkitSSHSession], Coroutine[Any, Any, None]],
+ *,
+ enable_cpr: bool,
+ ) -> None:
+ self.interact = interact
+ self.enable_cpr = enable_cpr
+ self.interact_task: asyncio.Task[None] | None = None
+ self._chan: Any | None = None
+ self.app_session: AppSession | None = None
+
+ # PipInput object, for sending input in the CLI.
+ # (This is something that we can use in the prompt_toolkit event loop,
+ # but still write date in manually.)
+ self._input: PipeInput | None = None
+ self._output: Vt100_Output | None = None
+
+ # Output object. Don't render to the real stdout, but write everything
+ # in the SSH channel.
+ class Stdout:
+ def write(s, data: str) -> None:
+ try:
+ if self._chan is not None:
+ self._chan.write(data.replace("\n", "\r\n"))
+ except BrokenPipeError:
+ pass # Channel not open for sending.
+
+ def isatty(s) -> bool:
+ return True
+
+ def flush(s) -> None:
+ pass
+
+ @property
+ def encoding(s) -> str:
+ assert self._chan is not None
+ return str(self._chan._orig_chan.get_encoding()[0])
+
+ self.stdout = cast(TextIO, Stdout())
+
+ def _get_size(self) -> Size:
+ """
+ Callable that returns the current `Size`, required by Vt100_Output.
+ """
+ if self._chan is None:
+ return Size(rows=20, columns=79)
+ else:
+ width, height, pixwidth, pixheight = self._chan.get_terminal_size()
+ return Size(rows=height, columns=width)
+
+ def connection_made(self, chan: Any) -> None:
+ self._chan = chan
+
+ def shell_requested(self) -> bool:
+ return True
+
+ def session_started(self) -> None:
+ self.interact_task = get_running_loop().create_task(self._interact())
+
+ async def _interact(self) -> None:
+ if self._chan is None:
+ # Should not happen.
+ raise Exception("`_interact` called before `connection_made`.")
+
+ if hasattr(self._chan, "set_line_mode") and self._chan._editor is not None:
+ # Disable the line editing provided by asyncssh. Prompt_toolkit
+ # provides the line editing.
+ self._chan.set_line_mode(False)
+
+ term = self._chan.get_terminal_type()
+
+ self._output = Vt100_Output(
+ self.stdout, self._get_size, term=term, enable_cpr=self.enable_cpr
+ )
+
+ with create_pipe_input() as self._input:
+ with create_app_session(input=self._input, output=self._output) as session:
+ self.app_session = session
+ try:
+ await self.interact(self)
+ except BaseException:
+ traceback.print_exc()
+ finally:
+ # Close the connection.
+ self._chan.close()
+ self._input.close()
+
+ def terminal_size_changed(
+ self, width: int, height: int, pixwidth: object, pixheight: object
+ ) -> None:
+ # Send resize event to the current application.
+ if self.app_session and self.app_session.app:
+ self.app_session.app._on_resize()
+
+ def data_received(self, data: str, datatype: object) -> None:
+ if self._input is None:
+ # Should not happen.
+ return
+
+ self._input.send_text(data)
+
+
+class PromptToolkitSSHServer(asyncssh.SSHServer):
+ """
+ Run a prompt_toolkit application over an asyncssh server.
+
+ This takes one argument, an `interact` function, which is called for each
+ connection. This should be an asynchronous function that runs the
+ prompt_toolkit applications. This function runs in an `AppSession`, which
+ means that we can have multiple UI interactions concurrently.
+
+ Example usage:
+
+ .. code:: python
+
+ async def interact(ssh_session: PromptToolkitSSHSession) -> None:
+ await yes_no_dialog("my title", "my text").run_async()
+
+ prompt_session = PromptSession()
+ text = await prompt_session.prompt_async("Type something: ")
+ print_formatted_text('You said: ', text)
+
+ server = PromptToolkitSSHServer(interact=interact)
+ loop = get_running_loop()
+ loop.run_until_complete(
+ asyncssh.create_server(
+ lambda: MySSHServer(interact),
+ "",
+ port,
+ server_host_keys=["/etc/ssh/..."],
+ )
+ )
+ loop.run_forever()
+
+ :param enable_cpr: When `True`, the default, try to detect whether the SSH
+ client runs in a terminal that responds to "cursor position requests".
+ That way, we can properly determine how much space there is available
+ for the UI (especially for drop down menus) to render.
+ """
+
+ def __init__(
+ self,
+ interact: Callable[[PromptToolkitSSHSession], Coroutine[Any, Any, None]],
+ *,
+ enable_cpr: bool = True,
+ ) -> None:
+ self.interact = interact
+ self.enable_cpr = enable_cpr
+
+ def begin_auth(self, username: str) -> bool:
+ # No authentication.
+ return False
+
+ def session_requested(self) -> PromptToolkitSSHSession:
+ return PromptToolkitSSHSession(self.interact, enable_cpr=self.enable_cpr)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..de902b43998562de06e92a586107381281922168
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__init__.py
@@ -0,0 +1,7 @@
+from __future__ import annotations
+
+from .server import TelnetServer
+
+__all__ = [
+ "TelnetServer",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..603a2a6f4a28f4a89427e131c7e3dd0ef07fcbe3
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/log.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/log.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..2fddc887765640c58f6058ce4c8e03e70268767c
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/log.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/protocol.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/protocol.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..edb59ecaceb20316e1267402380b4d38ecc1ec45
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/protocol.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/server.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/server.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..d5e9b45e86e5389d01eb676ff42b20a6393fe89f
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/__pycache__/server.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/log.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/log.py
new file mode 100644
index 0000000000000000000000000000000000000000..476dffc677489ca02d63ffbc2743f451cce44dc2
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/log.py
@@ -0,0 +1,13 @@
+"""
+Python logger for the telnet server.
+"""
+
+from __future__ import annotations
+
+import logging
+
+logger = logging.getLogger(__package__)
+
+__all__ = [
+ "logger",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/protocol.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/protocol.py
new file mode 100644
index 0000000000000000000000000000000000000000..58286e23e77b2dead0081c938fe813dcbd7aaa70
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/protocol.py
@@ -0,0 +1,209 @@
+"""
+Parser for the Telnet protocol. (Not a complete implementation of the telnet
+specification, but sufficient for a command line interface.)
+
+Inspired by `Twisted.conch.telnet`.
+"""
+
+from __future__ import annotations
+
+import struct
+from typing import Callable, Generator
+
+from .log import logger
+
+__all__ = [
+ "TelnetProtocolParser",
+]
+
+
+def int2byte(number: int) -> bytes:
+ return bytes((number,))
+
+
+# Telnet constants.
+NOP = int2byte(0)
+SGA = int2byte(3)
+
+IAC = int2byte(255)
+DO = int2byte(253)
+DONT = int2byte(254)
+LINEMODE = int2byte(34)
+SB = int2byte(250)
+WILL = int2byte(251)
+WONT = int2byte(252)
+MODE = int2byte(1)
+SE = int2byte(240)
+ECHO = int2byte(1)
+NAWS = int2byte(31)
+LINEMODE = int2byte(34)
+SUPPRESS_GO_AHEAD = int2byte(3)
+
+TTYPE = int2byte(24)
+SEND = int2byte(1)
+IS = int2byte(0)
+
+DM = int2byte(242)
+BRK = int2byte(243)
+IP = int2byte(244)
+AO = int2byte(245)
+AYT = int2byte(246)
+EC = int2byte(247)
+EL = int2byte(248)
+GA = int2byte(249)
+
+
+class TelnetProtocolParser:
+ """
+ Parser for the Telnet protocol.
+ Usage::
+
+ def data_received(data):
+ print(data)
+
+ def size_received(rows, columns):
+ print(rows, columns)
+
+ p = TelnetProtocolParser(data_received, size_received)
+ p.feed(binary_data)
+ """
+
+ def __init__(
+ self,
+ data_received_callback: Callable[[bytes], None],
+ size_received_callback: Callable[[int, int], None],
+ ttype_received_callback: Callable[[str], None],
+ ) -> None:
+ self.data_received_callback = data_received_callback
+ self.size_received_callback = size_received_callback
+ self.ttype_received_callback = ttype_received_callback
+
+ self._parser = self._parse_coroutine()
+ self._parser.send(None) # type: ignore
+
+ def received_data(self, data: bytes) -> None:
+ self.data_received_callback(data)
+
+ def do_received(self, data: bytes) -> None:
+ """Received telnet DO command."""
+ logger.info("DO %r", data)
+
+ def dont_received(self, data: bytes) -> None:
+ """Received telnet DONT command."""
+ logger.info("DONT %r", data)
+
+ def will_received(self, data: bytes) -> None:
+ """Received telnet WILL command."""
+ logger.info("WILL %r", data)
+
+ def wont_received(self, data: bytes) -> None:
+ """Received telnet WONT command."""
+ logger.info("WONT %r", data)
+
+ def command_received(self, command: bytes, data: bytes) -> None:
+ if command == DO:
+ self.do_received(data)
+
+ elif command == DONT:
+ self.dont_received(data)
+
+ elif command == WILL:
+ self.will_received(data)
+
+ elif command == WONT:
+ self.wont_received(data)
+
+ else:
+ logger.info("command received %r %r", command, data)
+
+ def naws(self, data: bytes) -> None:
+ """
+ Received NAWS. (Window dimensions.)
+ """
+ if len(data) == 4:
+ # NOTE: the first parameter of struct.unpack should be
+ # a 'str' object. Both on Py2/py3. This crashes on OSX
+ # otherwise.
+ columns, rows = struct.unpack("!HH", data)
+ self.size_received_callback(rows, columns)
+ else:
+ logger.warning("Wrong number of NAWS bytes")
+
+ def ttype(self, data: bytes) -> None:
+ """
+ Received terminal type.
+ """
+ subcmd, data = data[0:1], data[1:]
+ if subcmd == IS:
+ ttype = data.decode("ascii")
+ self.ttype_received_callback(ttype)
+ else:
+ logger.warning("Received a non-IS terminal type Subnegotiation")
+
+ def negotiate(self, data: bytes) -> None:
+ """
+ Got negotiate data.
+ """
+ command, payload = data[0:1], data[1:]
+
+ if command == NAWS:
+ self.naws(payload)
+ elif command == TTYPE:
+ self.ttype(payload)
+ else:
+ logger.info("Negotiate (%r got bytes)", len(data))
+
+ def _parse_coroutine(self) -> Generator[None, bytes, None]:
+ """
+ Parser state machine.
+ Every 'yield' expression returns the next byte.
+ """
+ while True:
+ d = yield
+
+ if d == int2byte(0):
+ pass # NOP
+
+ # Go to state escaped.
+ elif d == IAC:
+ d2 = yield
+
+ if d2 == IAC:
+ self.received_data(d2)
+
+ # Handle simple commands.
+ elif d2 in (NOP, DM, BRK, IP, AO, AYT, EC, EL, GA):
+ self.command_received(d2, b"")
+
+ # Handle IAC-[DO/DONT/WILL/WONT] commands.
+ elif d2 in (DO, DONT, WILL, WONT):
+ d3 = yield
+ self.command_received(d2, d3)
+
+ # Subnegotiation
+ elif d2 == SB:
+ # Consume everything until next IAC-SE
+ data = []
+
+ while True:
+ d3 = yield
+
+ if d3 == IAC:
+ d4 = yield
+ if d4 == SE:
+ break
+ else:
+ data.append(d4)
+ else:
+ data.append(d3)
+
+ self.negotiate(b"".join(data))
+ else:
+ self.received_data(d)
+
+ def feed(self, data: bytes) -> None:
+ """
+ Feed data to the parser.
+ """
+ for b in data:
+ self._parser.send(int2byte(b))
diff --git a/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/server.py b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/server.py
new file mode 100644
index 0000000000000000000000000000000000000000..bc2ad53064fd3a70e2533de75fea61f0af9c04b5
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/contrib/telnet/server.py
@@ -0,0 +1,428 @@
+"""
+Telnet server.
+"""
+
+from __future__ import annotations
+
+import asyncio
+import contextvars
+import socket
+from asyncio import get_running_loop
+from typing import Any, Callable, Coroutine, TextIO, cast
+
+from prompt_toolkit.application.current import create_app_session, get_app
+from prompt_toolkit.application.run_in_terminal import run_in_terminal
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.formatted_text import AnyFormattedText, to_formatted_text
+from prompt_toolkit.input import PipeInput, create_pipe_input
+from prompt_toolkit.output.vt100 import Vt100_Output
+from prompt_toolkit.renderer import print_formatted_text as print_formatted_text
+from prompt_toolkit.styles import BaseStyle, DummyStyle
+
+from .log import logger
+from .protocol import (
+ DO,
+ ECHO,
+ IAC,
+ LINEMODE,
+ MODE,
+ NAWS,
+ SB,
+ SE,
+ SEND,
+ SUPPRESS_GO_AHEAD,
+ TTYPE,
+ WILL,
+ TelnetProtocolParser,
+)
+
+__all__ = [
+ "TelnetServer",
+]
+
+
+def int2byte(number: int) -> bytes:
+ return bytes((number,))
+
+
+def _initialize_telnet(connection: socket.socket) -> None:
+ logger.info("Initializing telnet connection")
+
+ # Iac Do Linemode
+ connection.send(IAC + DO + LINEMODE)
+
+ # Suppress Go Ahead. (This seems important for Putty to do correct echoing.)
+ # This will allow bi-directional operation.
+ connection.send(IAC + WILL + SUPPRESS_GO_AHEAD)
+
+ # Iac sb
+ connection.send(IAC + SB + LINEMODE + MODE + int2byte(0) + IAC + SE)
+
+ # IAC Will Echo
+ connection.send(IAC + WILL + ECHO)
+
+ # Negotiate window size
+ connection.send(IAC + DO + NAWS)
+
+ # Negotiate terminal type
+ # Assume the client will accept the negotiation with `IAC + WILL + TTYPE`
+ connection.send(IAC + DO + TTYPE)
+
+ # We can then select the first terminal type supported by the client,
+ # which is generally the best type the client supports
+ # The client should reply with a `IAC + SB + TTYPE + IS + ttype + IAC + SE`
+ connection.send(IAC + SB + TTYPE + SEND + IAC + SE)
+
+
+class _ConnectionStdout:
+ """
+ Wrapper around socket which provides `write` and `flush` methods for the
+ Vt100_Output output.
+ """
+
+ def __init__(self, connection: socket.socket, encoding: str) -> None:
+ self._encoding = encoding
+ self._connection = connection
+ self._errors = "strict"
+ self._buffer: list[bytes] = []
+ self._closed = False
+
+ def write(self, data: str) -> None:
+ data = data.replace("\n", "\r\n")
+ self._buffer.append(data.encode(self._encoding, errors=self._errors))
+ self.flush()
+
+ def isatty(self) -> bool:
+ return True
+
+ def flush(self) -> None:
+ try:
+ if not self._closed:
+ self._connection.send(b"".join(self._buffer))
+ except OSError as e:
+ logger.warning(f"Couldn't send data over socket: {e}")
+
+ self._buffer = []
+
+ def close(self) -> None:
+ self._closed = True
+
+ @property
+ def encoding(self) -> str:
+ return self._encoding
+
+ @property
+ def errors(self) -> str:
+ return self._errors
+
+
+class TelnetConnection:
+ """
+ Class that represents one Telnet connection.
+ """
+
+ def __init__(
+ self,
+ conn: socket.socket,
+ addr: tuple[str, int],
+ interact: Callable[[TelnetConnection], Coroutine[Any, Any, None]],
+ server: TelnetServer,
+ encoding: str,
+ style: BaseStyle | None,
+ vt100_input: PipeInput,
+ enable_cpr: bool = True,
+ ) -> None:
+ self.conn = conn
+ self.addr = addr
+ self.interact = interact
+ self.server = server
+ self.encoding = encoding
+ self.style = style
+ self._closed = False
+ self._ready = asyncio.Event()
+ self.vt100_input = vt100_input
+ self.enable_cpr = enable_cpr
+ self.vt100_output: Vt100_Output | None = None
+
+ # Create "Output" object.
+ self.size = Size(rows=40, columns=79)
+
+ # Initialize.
+ _initialize_telnet(conn)
+
+ # Create output.
+ def get_size() -> Size:
+ return self.size
+
+ self.stdout = cast(TextIO, _ConnectionStdout(conn, encoding=encoding))
+
+ def data_received(data: bytes) -> None:
+ """TelnetProtocolParser 'data_received' callback"""
+ self.vt100_input.send_bytes(data)
+
+ def size_received(rows: int, columns: int) -> None:
+ """TelnetProtocolParser 'size_received' callback"""
+ self.size = Size(rows=rows, columns=columns)
+ if self.vt100_output is not None and self.context:
+ self.context.run(lambda: get_app()._on_resize())
+
+ def ttype_received(ttype: str) -> None:
+ """TelnetProtocolParser 'ttype_received' callback"""
+ self.vt100_output = Vt100_Output(
+ self.stdout, get_size, term=ttype, enable_cpr=enable_cpr
+ )
+ self._ready.set()
+
+ self.parser = TelnetProtocolParser(data_received, size_received, ttype_received)
+ self.context: contextvars.Context | None = None
+
+ async def run_application(self) -> None:
+ """
+ Run application.
+ """
+
+ def handle_incoming_data() -> None:
+ data = self.conn.recv(1024)
+ if data:
+ self.feed(data)
+ else:
+ # Connection closed by client.
+ logger.info("Connection closed by client. {!r} {!r}".format(*self.addr))
+ self.close()
+
+ # Add reader.
+ loop = get_running_loop()
+ loop.add_reader(self.conn, handle_incoming_data)
+
+ try:
+ # Wait for v100_output to be properly instantiated
+ await self._ready.wait()
+ with create_app_session(input=self.vt100_input, output=self.vt100_output):
+ self.context = contextvars.copy_context()
+ await self.interact(self)
+ finally:
+ self.close()
+
+ def feed(self, data: bytes) -> None:
+ """
+ Handler for incoming data. (Called by TelnetServer.)
+ """
+ self.parser.feed(data)
+
+ def close(self) -> None:
+ """
+ Closed by client.
+ """
+ if not self._closed:
+ self._closed = True
+
+ self.vt100_input.close()
+ get_running_loop().remove_reader(self.conn)
+ self.conn.close()
+ self.stdout.close()
+
+ def send(self, formatted_text: AnyFormattedText) -> None:
+ """
+ Send text to the client.
+ """
+ if self.vt100_output is None:
+ return
+ formatted_text = to_formatted_text(formatted_text)
+ print_formatted_text(
+ self.vt100_output, formatted_text, self.style or DummyStyle()
+ )
+
+ def send_above_prompt(self, formatted_text: AnyFormattedText) -> None:
+ """
+ Send text to the client.
+ This is asynchronous, returns a `Future`.
+ """
+ formatted_text = to_formatted_text(formatted_text)
+ return self._run_in_terminal(lambda: self.send(formatted_text))
+
+ def _run_in_terminal(self, func: Callable[[], None]) -> None:
+ # Make sure that when an application was active for this connection,
+ # that we print the text above the application.
+ if self.context:
+ self.context.run(run_in_terminal, func)
+ else:
+ raise RuntimeError("Called _run_in_terminal outside `run_application`.")
+
+ def erase_screen(self) -> None:
+ """
+ Erase the screen and move the cursor to the top.
+ """
+ if self.vt100_output is None:
+ return
+ self.vt100_output.erase_screen()
+ self.vt100_output.cursor_goto(0, 0)
+ self.vt100_output.flush()
+
+
+async def _dummy_interact(connection: TelnetConnection) -> None:
+ pass
+
+
+class TelnetServer:
+ """
+ Telnet server implementation.
+
+ Example::
+
+ async def interact(connection):
+ connection.send("Welcome")
+ session = PromptSession()
+ result = await session.prompt_async(message="Say something: ")
+ connection.send(f"You said: {result}\n")
+
+ async def main():
+ server = TelnetServer(interact=interact, port=2323)
+ await server.run()
+ """
+
+ def __init__(
+ self,
+ host: str = "127.0.0.1",
+ port: int = 23,
+ interact: Callable[
+ [TelnetConnection], Coroutine[Any, Any, None]
+ ] = _dummy_interact,
+ encoding: str = "utf-8",
+ style: BaseStyle | None = None,
+ enable_cpr: bool = True,
+ ) -> None:
+ self.host = host
+ self.port = port
+ self.interact = interact
+ self.encoding = encoding
+ self.style = style
+ self.enable_cpr = enable_cpr
+
+ self._run_task: asyncio.Task[None] | None = None
+ self._application_tasks: list[asyncio.Task[None]] = []
+
+ self.connections: set[TelnetConnection] = set()
+
+ @classmethod
+ def _create_socket(cls, host: str, port: int) -> socket.socket:
+ # Create and bind socket
+ s = socket.socket(socket.AF_INET, socket.SOCK_STREAM)
+ s.setsockopt(socket.SOL_SOCKET, socket.SO_REUSEADDR, 1)
+ s.bind((host, port))
+
+ s.listen(4)
+ return s
+
+ async def run(self, ready_cb: Callable[[], None] | None = None) -> None:
+ """
+ Run the telnet server, until this gets cancelled.
+
+ :param ready_cb: Callback that will be called at the point that we're
+ actually listening.
+ """
+ socket = self._create_socket(self.host, self.port)
+ logger.info(
+ "Listening for telnet connections on %s port %r", self.host, self.port
+ )
+
+ get_running_loop().add_reader(socket, lambda: self._accept(socket))
+
+ if ready_cb:
+ ready_cb()
+
+ try:
+ # Run forever, until cancelled.
+ await asyncio.Future()
+ finally:
+ get_running_loop().remove_reader(socket)
+ socket.close()
+
+ # Wait for all applications to finish.
+ for t in self._application_tasks:
+ t.cancel()
+
+ # (This is similar to
+ # `Application.cancel_and_wait_for_background_tasks`. We wait for the
+ # background tasks to complete, but don't propagate exceptions, because
+ # we can't use `ExceptionGroup` yet.)
+ if len(self._application_tasks) > 0:
+ await asyncio.wait(
+ self._application_tasks,
+ timeout=None,
+ return_when=asyncio.ALL_COMPLETED,
+ )
+
+ def start(self) -> None:
+ """
+ Deprecated: Use `.run()` instead.
+
+ Start the telnet server (stop by calling and awaiting `stop()`).
+ """
+ if self._run_task is not None:
+ # Already running.
+ return
+
+ self._run_task = get_running_loop().create_task(self.run())
+
+ async def stop(self) -> None:
+ """
+ Deprecated: Use `.run()` instead.
+
+ Stop a telnet server that was started using `.start()` and wait for the
+ cancellation to complete.
+ """
+ if self._run_task is not None:
+ self._run_task.cancel()
+ try:
+ await self._run_task
+ except asyncio.CancelledError:
+ pass
+
+ def _accept(self, listen_socket: socket.socket) -> None:
+ """
+ Accept new incoming connection.
+ """
+ conn, addr = listen_socket.accept()
+ logger.info("New connection %r %r", *addr)
+
+ # Run application for this connection.
+ async def run() -> None:
+ try:
+ with create_pipe_input() as vt100_input:
+ connection = TelnetConnection(
+ conn,
+ addr,
+ self.interact,
+ self,
+ encoding=self.encoding,
+ style=self.style,
+ vt100_input=vt100_input,
+ enable_cpr=self.enable_cpr,
+ )
+ self.connections.add(connection)
+
+ logger.info("Starting interaction %r %r", *addr)
+ try:
+ await connection.run_application()
+ finally:
+ self.connections.remove(connection)
+ logger.info("Stopping interaction %r %r", *addr)
+ except EOFError:
+ # Happens either when the connection is closed by the client
+ # (e.g., when the user types 'control-]', then 'quit' in the
+ # telnet client) or when the user types control-d in a prompt
+ # and this is not handled by the interact function.
+ logger.info("Unhandled EOFError in telnet application.")
+ except KeyboardInterrupt:
+ # Unhandled control-c propagated by a prompt.
+ logger.info("Unhandled KeyboardInterrupt in telnet application.")
+ except BaseException as e:
+ print(f"Got {type(e).__name__}", e)
+ import traceback
+
+ traceback.print_exc()
+ finally:
+ self._application_tasks.remove(task)
+
+ task = get_running_loop().create_task(run())
+ self._application_tasks.append(task)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/cursor_shapes.py b/lib/python3.12/site-packages/prompt_toolkit/cursor_shapes.py
new file mode 100644
index 0000000000000000000000000000000000000000..01d10926a26bb5932e556aa9bc4766b5682ef9cf
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/cursor_shapes.py
@@ -0,0 +1,117 @@
+from __future__ import annotations
+
+from abc import ABC, abstractmethod
+from enum import Enum
+from typing import TYPE_CHECKING, Any, Callable, Union
+
+from prompt_toolkit.enums import EditingMode
+from prompt_toolkit.key_binding.vi_state import InputMode
+
+if TYPE_CHECKING:
+ from .application import Application
+
+__all__ = [
+ "CursorShape",
+ "CursorShapeConfig",
+ "SimpleCursorShapeConfig",
+ "ModalCursorShapeConfig",
+ "DynamicCursorShapeConfig",
+ "to_cursor_shape_config",
+]
+
+
+class CursorShape(Enum):
+ # Default value that should tell the output implementation to never send
+ # cursor shape escape sequences. This is the default right now, because
+ # before this `CursorShape` functionality was introduced into
+ # prompt_toolkit itself, people had workarounds to send cursor shapes
+ # escapes into the terminal, by monkey patching some of prompt_toolkit's
+ # internals. We don't want the default prompt_toolkit implementation to
+ # interfere with that. E.g., IPython patches the `ViState.input_mode`
+ # property. See: https://github.com/ipython/ipython/pull/13501/files
+ _NEVER_CHANGE = "_NEVER_CHANGE"
+
+ BLOCK = "BLOCK"
+ BEAM = "BEAM"
+ UNDERLINE = "UNDERLINE"
+ BLINKING_BLOCK = "BLINKING_BLOCK"
+ BLINKING_BEAM = "BLINKING_BEAM"
+ BLINKING_UNDERLINE = "BLINKING_UNDERLINE"
+
+
+class CursorShapeConfig(ABC):
+ @abstractmethod
+ def get_cursor_shape(self, application: Application[Any]) -> CursorShape:
+ """
+ Return the cursor shape to be used in the current state.
+ """
+
+
+AnyCursorShapeConfig = Union[CursorShape, CursorShapeConfig, None]
+
+
+class SimpleCursorShapeConfig(CursorShapeConfig):
+ """
+ Always show the given cursor shape.
+ """
+
+ def __init__(self, cursor_shape: CursorShape = CursorShape._NEVER_CHANGE) -> None:
+ self.cursor_shape = cursor_shape
+
+ def get_cursor_shape(self, application: Application[Any]) -> CursorShape:
+ return self.cursor_shape
+
+
+class ModalCursorShapeConfig(CursorShapeConfig):
+ """
+ Show cursor shape according to the current input mode.
+ """
+
+ def get_cursor_shape(self, application: Application[Any]) -> CursorShape:
+ if application.editing_mode == EditingMode.VI:
+ if application.vi_state.input_mode in {
+ InputMode.NAVIGATION,
+ }:
+ return CursorShape.BLOCK
+ if application.vi_state.input_mode in {
+ InputMode.INSERT,
+ InputMode.INSERT_MULTIPLE,
+ }:
+ return CursorShape.BEAM
+ if application.vi_state.input_mode in {
+ InputMode.REPLACE,
+ InputMode.REPLACE_SINGLE,
+ }:
+ return CursorShape.UNDERLINE
+ elif application.editing_mode == EditingMode.EMACS:
+ # like vi's INSERT
+ return CursorShape.BEAM
+
+ # Default
+ return CursorShape.BLOCK
+
+
+class DynamicCursorShapeConfig(CursorShapeConfig):
+ def __init__(
+ self, get_cursor_shape_config: Callable[[], AnyCursorShapeConfig]
+ ) -> None:
+ self.get_cursor_shape_config = get_cursor_shape_config
+
+ def get_cursor_shape(self, application: Application[Any]) -> CursorShape:
+ return to_cursor_shape_config(self.get_cursor_shape_config()).get_cursor_shape(
+ application
+ )
+
+
+def to_cursor_shape_config(value: AnyCursorShapeConfig) -> CursorShapeConfig:
+ """
+ Take a `CursorShape` instance or `CursorShapeConfig` and turn it into a
+ `CursorShapeConfig`.
+ """
+ if value is None:
+ return SimpleCursorShapeConfig()
+
+ if isinstance(value, CursorShape):
+ return SimpleCursorShapeConfig(value)
+
+ return value
diff --git a/lib/python3.12/site-packages/prompt_toolkit/data_structures.py b/lib/python3.12/site-packages/prompt_toolkit/data_structures.py
new file mode 100644
index 0000000000000000000000000000000000000000..27dd4585438eab2a684375604db6a765617b03b4
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/data_structures.py
@@ -0,0 +1,18 @@
+from __future__ import annotations
+
+from typing import NamedTuple
+
+__all__ = [
+ "Point",
+ "Size",
+]
+
+
+class Point(NamedTuple):
+ x: int
+ y: int
+
+
+class Size(NamedTuple):
+ rows: int
+ columns: int
diff --git a/lib/python3.12/site-packages/prompt_toolkit/document.py b/lib/python3.12/site-packages/prompt_toolkit/document.py
new file mode 100644
index 0000000000000000000000000000000000000000..c9fbd9a369671fd73d197552c4ff8ac7eb69ceab
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/document.py
@@ -0,0 +1,1182 @@
+"""
+The `Document` that implements all the text operations/querying.
+"""
+
+from __future__ import annotations
+
+import bisect
+import re
+import string
+import weakref
+from typing import Callable, Dict, Iterable, List, NoReturn, Pattern, cast
+
+from .clipboard import ClipboardData
+from .filters import vi_mode
+from .selection import PasteMode, SelectionState, SelectionType
+
+__all__ = [
+ "Document",
+]
+
+
+# Regex for finding "words" in documents. (We consider a group of alnum
+# characters a word, but also a group of special characters a word, as long as
+# it doesn't contain a space.)
+# (This is a 'word' in Vi.)
+_FIND_WORD_RE = re.compile(r"([a-zA-Z0-9_]+|[^a-zA-Z0-9_\s]+)")
+_FIND_CURRENT_WORD_RE = re.compile(r"^([a-zA-Z0-9_]+|[^a-zA-Z0-9_\s]+)")
+_FIND_CURRENT_WORD_INCLUDE_TRAILING_WHITESPACE_RE = re.compile(
+ r"^(([a-zA-Z0-9_]+|[^a-zA-Z0-9_\s]+)\s*)"
+)
+
+# Regex for finding "WORDS" in documents.
+# (This is a 'WORD in Vi.)
+_FIND_BIG_WORD_RE = re.compile(r"([^\s]+)")
+_FIND_CURRENT_BIG_WORD_RE = re.compile(r"^([^\s]+)")
+_FIND_CURRENT_BIG_WORD_INCLUDE_TRAILING_WHITESPACE_RE = re.compile(r"^([^\s]+\s*)")
+
+# Share the Document._cache between all Document instances.
+# (Document instances are considered immutable. That means that if another
+# `Document` is constructed with the same text, it should have the same
+# `_DocumentCache`.)
+_text_to_document_cache: dict[str, _DocumentCache] = cast(
+ Dict[str, "_DocumentCache"],
+ weakref.WeakValueDictionary(), # Maps document.text to DocumentCache instance.
+)
+
+
+class _ImmutableLineList(List[str]):
+ """
+ Some protection for our 'lines' list, which is assumed to be immutable in the cache.
+ (Useful for detecting obvious bugs.)
+ """
+
+ def _error(self, *a: object, **kw: object) -> NoReturn:
+ raise NotImplementedError("Attempt to modify an immutable list.")
+
+ __setitem__ = _error
+ append = _error
+ clear = _error
+ extend = _error
+ insert = _error
+ pop = _error
+ remove = _error
+ reverse = _error
+ sort = _error
+
+
+class _DocumentCache:
+ def __init__(self) -> None:
+ #: List of lines for the Document text.
+ self.lines: _ImmutableLineList | None = None
+
+ #: List of index positions, pointing to the start of all the lines.
+ self.line_indexes: list[int] | None = None
+
+
+class Document:
+ """
+ This is a immutable class around the text and cursor position, and contains
+ methods for querying this data, e.g. to give the text before the cursor.
+
+ This class is usually instantiated by a :class:`~prompt_toolkit.buffer.Buffer`
+ object, and accessed as the `document` property of that class.
+
+ :param text: string
+ :param cursor_position: int
+ :param selection: :class:`.SelectionState`
+ """
+
+ __slots__ = ("_text", "_cursor_position", "_selection", "_cache")
+
+ def __init__(
+ self,
+ text: str = "",
+ cursor_position: int | None = None,
+ selection: SelectionState | None = None,
+ ) -> None:
+ # Check cursor position. It can also be right after the end. (Where we
+ # insert text.)
+ assert cursor_position is None or cursor_position <= len(text), AssertionError(
+ f"cursor_position={cursor_position!r}, len_text={len(text)!r}"
+ )
+
+ # By default, if no cursor position was given, make sure to put the
+ # cursor position is at the end of the document. This is what makes
+ # sense in most places.
+ if cursor_position is None:
+ cursor_position = len(text)
+
+ # Keep these attributes private. A `Document` really has to be
+ # considered to be immutable, because otherwise the caching will break
+ # things. Because of that, we wrap these into read-only properties.
+ self._text = text
+ self._cursor_position = cursor_position
+ self._selection = selection
+
+ # Cache for lines/indexes. (Shared with other Document instances that
+ # contain the same text.
+ try:
+ self._cache = _text_to_document_cache[self.text]
+ except KeyError:
+ self._cache = _DocumentCache()
+ _text_to_document_cache[self.text] = self._cache
+
+ # XX: For some reason, above, we can't use 'WeakValueDictionary.setdefault'.
+ # This fails in Pypy3. `self._cache` becomes None, because that's what
+ # 'setdefault' returns.
+ # self._cache = _text_to_document_cache.setdefault(self.text, _DocumentCache())
+ # assert self._cache
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.text!r}, {self.cursor_position!r})"
+
+ def __eq__(self, other: object) -> bool:
+ if not isinstance(other, Document):
+ return False
+
+ return (
+ self.text == other.text
+ and self.cursor_position == other.cursor_position
+ and self.selection == other.selection
+ )
+
+ @property
+ def text(self) -> str:
+ "The document text."
+ return self._text
+
+ @property
+ def cursor_position(self) -> int:
+ "The document cursor position."
+ return self._cursor_position
+
+ @property
+ def selection(self) -> SelectionState | None:
+ ":class:`.SelectionState` object."
+ return self._selection
+
+ @property
+ def current_char(self) -> str:
+ """Return character under cursor or an empty string."""
+ return self._get_char_relative_to_cursor(0) or ""
+
+ @property
+ def char_before_cursor(self) -> str:
+ """Return character before the cursor or an empty string."""
+ return self._get_char_relative_to_cursor(-1) or ""
+
+ @property
+ def text_before_cursor(self) -> str:
+ return self.text[: self.cursor_position :]
+
+ @property
+ def text_after_cursor(self) -> str:
+ return self.text[self.cursor_position :]
+
+ @property
+ def current_line_before_cursor(self) -> str:
+ """Text from the start of the line until the cursor."""
+ _, _, text = self.text_before_cursor.rpartition("\n")
+ return text
+
+ @property
+ def current_line_after_cursor(self) -> str:
+ """Text from the cursor until the end of the line."""
+ text, _, _ = self.text_after_cursor.partition("\n")
+ return text
+
+ @property
+ def lines(self) -> list[str]:
+ """
+ Array of all the lines.
+ """
+ # Cache, because this one is reused very often.
+ if self._cache.lines is None:
+ self._cache.lines = _ImmutableLineList(self.text.split("\n"))
+
+ return self._cache.lines
+
+ @property
+ def _line_start_indexes(self) -> list[int]:
+ """
+ Array pointing to the start indexes of all the lines.
+ """
+ # Cache, because this is often reused. (If it is used, it's often used
+ # many times. And this has to be fast for editing big documents!)
+ if self._cache.line_indexes is None:
+ # Create list of line lengths.
+ line_lengths = map(len, self.lines)
+
+ # Calculate cumulative sums.
+ indexes = [0]
+ append = indexes.append
+ pos = 0
+
+ for line_length in line_lengths:
+ pos += line_length + 1
+ append(pos)
+
+ # Remove the last item. (This is not a new line.)
+ if len(indexes) > 1:
+ indexes.pop()
+
+ self._cache.line_indexes = indexes
+
+ return self._cache.line_indexes
+
+ @property
+ def lines_from_current(self) -> list[str]:
+ """
+ Array of the lines starting from the current line, until the last line.
+ """
+ return self.lines[self.cursor_position_row :]
+
+ @property
+ def line_count(self) -> int:
+ r"""Return the number of lines in this document. If the document ends
+ with a trailing \n, that counts as the beginning of a new line."""
+ return len(self.lines)
+
+ @property
+ def current_line(self) -> str:
+ """Return the text on the line where the cursor is. (when the input
+ consists of just one line, it equals `text`."""
+ return self.current_line_before_cursor + self.current_line_after_cursor
+
+ @property
+ def leading_whitespace_in_current_line(self) -> str:
+ """The leading whitespace in the left margin of the current line."""
+ current_line = self.current_line
+ length = len(current_line) - len(current_line.lstrip())
+ return current_line[:length]
+
+ def _get_char_relative_to_cursor(self, offset: int = 0) -> str:
+ """
+ Return character relative to cursor position, or empty string
+ """
+ try:
+ return self.text[self.cursor_position + offset]
+ except IndexError:
+ return ""
+
+ @property
+ def on_first_line(self) -> bool:
+ """
+ True when we are at the first line.
+ """
+ return self.cursor_position_row == 0
+
+ @property
+ def on_last_line(self) -> bool:
+ """
+ True when we are at the last line.
+ """
+ return self.cursor_position_row == self.line_count - 1
+
+ @property
+ def cursor_position_row(self) -> int:
+ """
+ Current row. (0-based.)
+ """
+ row, _ = self._find_line_start_index(self.cursor_position)
+ return row
+
+ @property
+ def cursor_position_col(self) -> int:
+ """
+ Current column. (0-based.)
+ """
+ # (Don't use self.text_before_cursor to calculate this. Creating
+ # substrings and doing rsplit is too expensive for getting the cursor
+ # position.)
+ _, line_start_index = self._find_line_start_index(self.cursor_position)
+ return self.cursor_position - line_start_index
+
+ def _find_line_start_index(self, index: int) -> tuple[int, int]:
+ """
+ For the index of a character at a certain line, calculate the index of
+ the first character on that line.
+
+ Return (row, index) tuple.
+ """
+ indexes = self._line_start_indexes
+
+ pos = bisect.bisect_right(indexes, index) - 1
+ return pos, indexes[pos]
+
+ def translate_index_to_position(self, index: int) -> tuple[int, int]:
+ """
+ Given an index for the text, return the corresponding (row, col) tuple.
+ (0-based. Returns (0, 0) for index=0.)
+ """
+ # Find start of this line.
+ row, row_index = self._find_line_start_index(index)
+ col = index - row_index
+
+ return row, col
+
+ def translate_row_col_to_index(self, row: int, col: int) -> int:
+ """
+ Given a (row, col) tuple, return the corresponding index.
+ (Row and col params are 0-based.)
+
+ Negative row/col values are turned into zero.
+ """
+ try:
+ result = self._line_start_indexes[row]
+ line = self.lines[row]
+ except IndexError:
+ if row < 0:
+ result = self._line_start_indexes[0]
+ line = self.lines[0]
+ else:
+ result = self._line_start_indexes[-1]
+ line = self.lines[-1]
+
+ result += max(0, min(col, len(line)))
+
+ # Keep in range. (len(self.text) is included, because the cursor can be
+ # right after the end of the text as well.)
+ result = max(0, min(result, len(self.text)))
+ return result
+
+ @property
+ def is_cursor_at_the_end(self) -> bool:
+ """True when the cursor is at the end of the text."""
+ return self.cursor_position == len(self.text)
+
+ @property
+ def is_cursor_at_the_end_of_line(self) -> bool:
+ """True when the cursor is at the end of this line."""
+ return self.current_char in ("\n", "")
+
+ def has_match_at_current_position(self, sub: str) -> bool:
+ """
+ `True` when this substring is found at the cursor position.
+ """
+ return self.text.find(sub, self.cursor_position) == self.cursor_position
+
+ def find(
+ self,
+ sub: str,
+ in_current_line: bool = False,
+ include_current_position: bool = False,
+ ignore_case: bool = False,
+ count: int = 1,
+ ) -> int | None:
+ """
+ Find `text` after the cursor, return position relative to the cursor
+ position. Return `None` if nothing was found.
+
+ :param count: Find the n-th occurrence.
+ """
+ assert isinstance(ignore_case, bool)
+
+ if in_current_line:
+ text = self.current_line_after_cursor
+ else:
+ text = self.text_after_cursor
+
+ if not include_current_position:
+ if len(text) == 0:
+ return None # (Otherwise, we always get a match for the empty string.)
+ else:
+ text = text[1:]
+
+ flags = re.IGNORECASE if ignore_case else 0
+ iterator = re.finditer(re.escape(sub), text, flags)
+
+ try:
+ for i, match in enumerate(iterator):
+ if i + 1 == count:
+ if include_current_position:
+ return match.start(0)
+ else:
+ return match.start(0) + 1
+ except StopIteration:
+ pass
+ return None
+
+ def find_all(self, sub: str, ignore_case: bool = False) -> list[int]:
+ """
+ Find all occurrences of the substring. Return a list of absolute
+ positions in the document.
+ """
+ flags = re.IGNORECASE if ignore_case else 0
+ return [a.start() for a in re.finditer(re.escape(sub), self.text, flags)]
+
+ def find_backwards(
+ self,
+ sub: str,
+ in_current_line: bool = False,
+ ignore_case: bool = False,
+ count: int = 1,
+ ) -> int | None:
+ """
+ Find `text` before the cursor, return position relative to the cursor
+ position. Return `None` if nothing was found.
+
+ :param count: Find the n-th occurrence.
+ """
+ if in_current_line:
+ before_cursor = self.current_line_before_cursor[::-1]
+ else:
+ before_cursor = self.text_before_cursor[::-1]
+
+ flags = re.IGNORECASE if ignore_case else 0
+ iterator = re.finditer(re.escape(sub[::-1]), before_cursor, flags)
+
+ try:
+ for i, match in enumerate(iterator):
+ if i + 1 == count:
+ return -match.start(0) - len(sub)
+ except StopIteration:
+ pass
+ return None
+
+ def get_word_before_cursor(
+ self, WORD: bool = False, pattern: Pattern[str] | None = None
+ ) -> str:
+ """
+ Give the word before the cursor.
+ If we have whitespace before the cursor this returns an empty string.
+
+ :param pattern: (None or compiled regex). When given, use this regex
+ pattern.
+ """
+ if self._is_word_before_cursor_complete(WORD=WORD, pattern=pattern):
+ # Space before the cursor or no text before cursor.
+ return ""
+
+ text_before_cursor = self.text_before_cursor
+ start = self.find_start_of_previous_word(WORD=WORD, pattern=pattern) or 0
+
+ return text_before_cursor[len(text_before_cursor) + start :]
+
+ def _is_word_before_cursor_complete(
+ self, WORD: bool = False, pattern: Pattern[str] | None = None
+ ) -> bool:
+ if pattern:
+ return self.find_start_of_previous_word(WORD=WORD, pattern=pattern) is None
+ else:
+ return (
+ self.text_before_cursor == "" or self.text_before_cursor[-1:].isspace()
+ )
+
+ def find_start_of_previous_word(
+ self, count: int = 1, WORD: bool = False, pattern: Pattern[str] | None = None
+ ) -> int | None:
+ """
+ Return an index relative to the cursor position pointing to the start
+ of the previous word. Return `None` if nothing was found.
+
+ :param pattern: (None or compiled regex). When given, use this regex
+ pattern.
+ """
+ assert not (WORD and pattern)
+
+ # Reverse the text before the cursor, in order to do an efficient
+ # backwards search.
+ text_before_cursor = self.text_before_cursor[::-1]
+
+ if pattern:
+ regex = pattern
+ elif WORD:
+ regex = _FIND_BIG_WORD_RE
+ else:
+ regex = _FIND_WORD_RE
+
+ iterator = regex.finditer(text_before_cursor)
+
+ try:
+ for i, match in enumerate(iterator):
+ if i + 1 == count:
+ return -match.end(0)
+ except StopIteration:
+ pass
+ return None
+
+ def find_boundaries_of_current_word(
+ self,
+ WORD: bool = False,
+ include_leading_whitespace: bool = False,
+ include_trailing_whitespace: bool = False,
+ ) -> tuple[int, int]:
+ """
+ Return the relative boundaries (startpos, endpos) of the current word under the
+ cursor. (This is at the current line, because line boundaries obviously
+ don't belong to any word.)
+ If not on a word, this returns (0,0)
+ """
+ text_before_cursor = self.current_line_before_cursor[::-1]
+ text_after_cursor = self.current_line_after_cursor
+
+ def get_regex(include_whitespace: bool) -> Pattern[str]:
+ return {
+ (False, False): _FIND_CURRENT_WORD_RE,
+ (False, True): _FIND_CURRENT_WORD_INCLUDE_TRAILING_WHITESPACE_RE,
+ (True, False): _FIND_CURRENT_BIG_WORD_RE,
+ (True, True): _FIND_CURRENT_BIG_WORD_INCLUDE_TRAILING_WHITESPACE_RE,
+ }[(WORD, include_whitespace)]
+
+ match_before = get_regex(include_leading_whitespace).search(text_before_cursor)
+ match_after = get_regex(include_trailing_whitespace).search(text_after_cursor)
+
+ # When there is a match before and after, and we're not looking for
+ # WORDs, make sure that both the part before and after the cursor are
+ # either in the [a-zA-Z_] alphabet or not. Otherwise, drop the part
+ # before the cursor.
+ if not WORD and match_before and match_after:
+ c1 = self.text[self.cursor_position - 1]
+ c2 = self.text[self.cursor_position]
+ alphabet = string.ascii_letters + "0123456789_"
+
+ if (c1 in alphabet) != (c2 in alphabet):
+ match_before = None
+
+ return (
+ -match_before.end(1) if match_before else 0,
+ match_after.end(1) if match_after else 0,
+ )
+
+ def get_word_under_cursor(self, WORD: bool = False) -> str:
+ """
+ Return the word, currently below the cursor.
+ This returns an empty string when the cursor is on a whitespace region.
+ """
+ start, end = self.find_boundaries_of_current_word(WORD=WORD)
+ return self.text[self.cursor_position + start : self.cursor_position + end]
+
+ def find_next_word_beginning(
+ self, count: int = 1, WORD: bool = False
+ ) -> int | None:
+ """
+ Return an index relative to the cursor position pointing to the start
+ of the next word. Return `None` if nothing was found.
+ """
+ if count < 0:
+ return self.find_previous_word_beginning(count=-count, WORD=WORD)
+
+ regex = _FIND_BIG_WORD_RE if WORD else _FIND_WORD_RE
+ iterator = regex.finditer(self.text_after_cursor)
+
+ try:
+ for i, match in enumerate(iterator):
+ # Take first match, unless it's the word on which we're right now.
+ if i == 0 and match.start(1) == 0:
+ count += 1
+
+ if i + 1 == count:
+ return match.start(1)
+ except StopIteration:
+ pass
+ return None
+
+ def find_next_word_ending(
+ self, include_current_position: bool = False, count: int = 1, WORD: bool = False
+ ) -> int | None:
+ """
+ Return an index relative to the cursor position pointing to the end
+ of the next word. Return `None` if nothing was found.
+ """
+ if count < 0:
+ return self.find_previous_word_ending(count=-count, WORD=WORD)
+
+ if include_current_position:
+ text = self.text_after_cursor
+ else:
+ text = self.text_after_cursor[1:]
+
+ regex = _FIND_BIG_WORD_RE if WORD else _FIND_WORD_RE
+ iterable = regex.finditer(text)
+
+ try:
+ for i, match in enumerate(iterable):
+ if i + 1 == count:
+ value = match.end(1)
+
+ if include_current_position:
+ return value
+ else:
+ return value + 1
+
+ except StopIteration:
+ pass
+ return None
+
+ def find_previous_word_beginning(
+ self, count: int = 1, WORD: bool = False
+ ) -> int | None:
+ """
+ Return an index relative to the cursor position pointing to the start
+ of the previous word. Return `None` if nothing was found.
+ """
+ if count < 0:
+ return self.find_next_word_beginning(count=-count, WORD=WORD)
+
+ regex = _FIND_BIG_WORD_RE if WORD else _FIND_WORD_RE
+ iterator = regex.finditer(self.text_before_cursor[::-1])
+
+ try:
+ for i, match in enumerate(iterator):
+ if i + 1 == count:
+ return -match.end(1)
+ except StopIteration:
+ pass
+ return None
+
+ def find_previous_word_ending(
+ self, count: int = 1, WORD: bool = False
+ ) -> int | None:
+ """
+ Return an index relative to the cursor position pointing to the end
+ of the previous word. Return `None` if nothing was found.
+ """
+ if count < 0:
+ return self.find_next_word_ending(count=-count, WORD=WORD)
+
+ text_before_cursor = self.text_after_cursor[:1] + self.text_before_cursor[::-1]
+
+ regex = _FIND_BIG_WORD_RE if WORD else _FIND_WORD_RE
+ iterator = regex.finditer(text_before_cursor)
+
+ try:
+ for i, match in enumerate(iterator):
+ # Take first match, unless it's the word on which we're right now.
+ if i == 0 and match.start(1) == 0:
+ count += 1
+
+ if i + 1 == count:
+ return -match.start(1) + 1
+ except StopIteration:
+ pass
+ return None
+
+ def find_next_matching_line(
+ self, match_func: Callable[[str], bool], count: int = 1
+ ) -> int | None:
+ """
+ Look downwards for empty lines.
+ Return the line index, relative to the current line.
+ """
+ result = None
+
+ for index, line in enumerate(self.lines[self.cursor_position_row + 1 :]):
+ if match_func(line):
+ result = 1 + index
+ count -= 1
+
+ if count == 0:
+ break
+
+ return result
+
+ def find_previous_matching_line(
+ self, match_func: Callable[[str], bool], count: int = 1
+ ) -> int | None:
+ """
+ Look upwards for empty lines.
+ Return the line index, relative to the current line.
+ """
+ result = None
+
+ for index, line in enumerate(self.lines[: self.cursor_position_row][::-1]):
+ if match_func(line):
+ result = -1 - index
+ count -= 1
+
+ if count == 0:
+ break
+
+ return result
+
+ def get_cursor_left_position(self, count: int = 1) -> int:
+ """
+ Relative position for cursor left.
+ """
+ if count < 0:
+ return self.get_cursor_right_position(-count)
+
+ return -min(self.cursor_position_col, count)
+
+ def get_cursor_right_position(self, count: int = 1) -> int:
+ """
+ Relative position for cursor_right.
+ """
+ if count < 0:
+ return self.get_cursor_left_position(-count)
+
+ return min(count, len(self.current_line_after_cursor))
+
+ def get_cursor_up_position(
+ self, count: int = 1, preferred_column: int | None = None
+ ) -> int:
+ """
+ Return the relative cursor position (character index) where we would be if the
+ user pressed the arrow-up button.
+
+ :param preferred_column: When given, go to this column instead of
+ staying at the current column.
+ """
+ assert count >= 1
+ column = (
+ self.cursor_position_col if preferred_column is None else preferred_column
+ )
+
+ return (
+ self.translate_row_col_to_index(
+ max(0, self.cursor_position_row - count), column
+ )
+ - self.cursor_position
+ )
+
+ def get_cursor_down_position(
+ self, count: int = 1, preferred_column: int | None = None
+ ) -> int:
+ """
+ Return the relative cursor position (character index) where we would be if the
+ user pressed the arrow-down button.
+
+ :param preferred_column: When given, go to this column instead of
+ staying at the current column.
+ """
+ assert count >= 1
+ column = (
+ self.cursor_position_col if preferred_column is None else preferred_column
+ )
+
+ return (
+ self.translate_row_col_to_index(self.cursor_position_row + count, column)
+ - self.cursor_position
+ )
+
+ def find_enclosing_bracket_right(
+ self, left_ch: str, right_ch: str, end_pos: int | None = None
+ ) -> int | None:
+ """
+ Find the right bracket enclosing current position. Return the relative
+ position to the cursor position.
+
+ When `end_pos` is given, don't look past the position.
+ """
+ if self.current_char == right_ch:
+ return 0
+
+ if end_pos is None:
+ end_pos = len(self.text)
+ else:
+ end_pos = min(len(self.text), end_pos)
+
+ stack = 1
+
+ # Look forward.
+ for i in range(self.cursor_position + 1, end_pos):
+ c = self.text[i]
+
+ if c == left_ch:
+ stack += 1
+ elif c == right_ch:
+ stack -= 1
+
+ if stack == 0:
+ return i - self.cursor_position
+
+ return None
+
+ def find_enclosing_bracket_left(
+ self, left_ch: str, right_ch: str, start_pos: int | None = None
+ ) -> int | None:
+ """
+ Find the left bracket enclosing current position. Return the relative
+ position to the cursor position.
+
+ When `start_pos` is given, don't look past the position.
+ """
+ if self.current_char == left_ch:
+ return 0
+
+ if start_pos is None:
+ start_pos = 0
+ else:
+ start_pos = max(0, start_pos)
+
+ stack = 1
+
+ # Look backward.
+ for i in range(self.cursor_position - 1, start_pos - 1, -1):
+ c = self.text[i]
+
+ if c == right_ch:
+ stack += 1
+ elif c == left_ch:
+ stack -= 1
+
+ if stack == 0:
+ return i - self.cursor_position
+
+ return None
+
+ def find_matching_bracket_position(
+ self, start_pos: int | None = None, end_pos: int | None = None
+ ) -> int:
+ """
+ Return relative cursor position of matching [, (, { or < bracket.
+
+ When `start_pos` or `end_pos` are given. Don't look past the positions.
+ """
+
+ # Look for a match.
+ for pair in "()", "[]", "{}", "<>":
+ A = pair[0]
+ B = pair[1]
+ if self.current_char == A:
+ return self.find_enclosing_bracket_right(A, B, end_pos=end_pos) or 0
+ elif self.current_char == B:
+ return self.find_enclosing_bracket_left(A, B, start_pos=start_pos) or 0
+
+ return 0
+
+ def get_start_of_document_position(self) -> int:
+ """Relative position for the start of the document."""
+ return -self.cursor_position
+
+ def get_end_of_document_position(self) -> int:
+ """Relative position for the end of the document."""
+ return len(self.text) - self.cursor_position
+
+ def get_start_of_line_position(self, after_whitespace: bool = False) -> int:
+ """Relative position for the start of this line."""
+ if after_whitespace:
+ current_line = self.current_line
+ return (
+ len(current_line)
+ - len(current_line.lstrip())
+ - self.cursor_position_col
+ )
+ else:
+ return -len(self.current_line_before_cursor)
+
+ def get_end_of_line_position(self) -> int:
+ """Relative position for the end of this line."""
+ return len(self.current_line_after_cursor)
+
+ def last_non_blank_of_current_line_position(self) -> int:
+ """
+ Relative position for the last non blank character of this line.
+ """
+ return len(self.current_line.rstrip()) - self.cursor_position_col - 1
+
+ def get_column_cursor_position(self, column: int) -> int:
+ """
+ Return the relative cursor position for this column at the current
+ line. (It will stay between the boundaries of the line in case of a
+ larger number.)
+ """
+ line_length = len(self.current_line)
+ current_column = self.cursor_position_col
+ column = max(0, min(line_length, column))
+
+ return column - current_column
+
+ def selection_range(
+ self,
+ ) -> tuple[
+ int, int
+ ]: # XXX: shouldn't this return `None` if there is no selection???
+ """
+ Return (from, to) tuple of the selection.
+ start and end position are included.
+
+ This doesn't take the selection type into account. Use
+ `selection_ranges` instead.
+ """
+ if self.selection:
+ from_, to = sorted(
+ [self.cursor_position, self.selection.original_cursor_position]
+ )
+ else:
+ from_, to = self.cursor_position, self.cursor_position
+
+ return from_, to
+
+ def selection_ranges(self) -> Iterable[tuple[int, int]]:
+ """
+ Return a list of `(from, to)` tuples for the selection or none if
+ nothing was selected. The upper boundary is not included.
+
+ This will yield several (from, to) tuples in case of a BLOCK selection.
+ This will return zero ranges, like (8,8) for empty lines in a block
+ selection.
+ """
+ if self.selection:
+ from_, to = sorted(
+ [self.cursor_position, self.selection.original_cursor_position]
+ )
+
+ if self.selection.type == SelectionType.BLOCK:
+ from_line, from_column = self.translate_index_to_position(from_)
+ to_line, to_column = self.translate_index_to_position(to)
+ from_column, to_column = sorted([from_column, to_column])
+ lines = self.lines
+
+ if vi_mode():
+ to_column += 1
+
+ for l in range(from_line, to_line + 1):
+ line_length = len(lines[l])
+
+ if from_column <= line_length:
+ yield (
+ self.translate_row_col_to_index(l, from_column),
+ self.translate_row_col_to_index(
+ l, min(line_length, to_column)
+ ),
+ )
+ else:
+ # In case of a LINES selection, go to the start/end of the lines.
+ if self.selection.type == SelectionType.LINES:
+ from_ = max(0, self.text.rfind("\n", 0, from_) + 1)
+
+ if self.text.find("\n", to) >= 0:
+ to = self.text.find("\n", to)
+ else:
+ to = len(self.text) - 1
+
+ # In Vi mode, the upper boundary is always included. For Emacs,
+ # that's not the case.
+ if vi_mode():
+ to += 1
+
+ yield from_, to
+
+ def selection_range_at_line(self, row: int) -> tuple[int, int] | None:
+ """
+ If the selection spans a portion of the given line, return a (from, to) tuple.
+
+ The returned upper boundary is not included in the selection, so
+ `(0, 0)` is an empty selection. `(0, 1)`, is a one character selection.
+
+ Returns None if the selection doesn't cover this line at all.
+ """
+ if self.selection:
+ line = self.lines[row]
+
+ row_start = self.translate_row_col_to_index(row, 0)
+ row_end = self.translate_row_col_to_index(row, len(line))
+
+ from_, to = sorted(
+ [self.cursor_position, self.selection.original_cursor_position]
+ )
+
+ # Take the intersection of the current line and the selection.
+ intersection_start = max(row_start, from_)
+ intersection_end = min(row_end, to)
+
+ if intersection_start <= intersection_end:
+ if self.selection.type == SelectionType.LINES:
+ intersection_start = row_start
+ intersection_end = row_end
+
+ elif self.selection.type == SelectionType.BLOCK:
+ _, col1 = self.translate_index_to_position(from_)
+ _, col2 = self.translate_index_to_position(to)
+ col1, col2 = sorted([col1, col2])
+
+ if col1 > len(line):
+ return None # Block selection doesn't cross this line.
+
+ intersection_start = self.translate_row_col_to_index(row, col1)
+ intersection_end = self.translate_row_col_to_index(row, col2)
+
+ _, from_column = self.translate_index_to_position(intersection_start)
+ _, to_column = self.translate_index_to_position(intersection_end)
+
+ # In Vi mode, the upper boundary is always included. For Emacs
+ # mode, that's not the case.
+ if vi_mode():
+ to_column += 1
+
+ return from_column, to_column
+ return None
+
+ def cut_selection(self) -> tuple[Document, ClipboardData]:
+ """
+ Return a (:class:`.Document`, :class:`.ClipboardData`) tuple, where the
+ document represents the new document when the selection is cut, and the
+ clipboard data, represents whatever has to be put on the clipboard.
+ """
+ if self.selection:
+ cut_parts = []
+ remaining_parts = []
+ new_cursor_position = self.cursor_position
+
+ last_to = 0
+ for from_, to in self.selection_ranges():
+ if last_to == 0:
+ new_cursor_position = from_
+
+ remaining_parts.append(self.text[last_to:from_])
+ cut_parts.append(self.text[from_:to])
+ last_to = to
+
+ remaining_parts.append(self.text[last_to:])
+
+ cut_text = "\n".join(cut_parts)
+ remaining_text = "".join(remaining_parts)
+
+ # In case of a LINES selection, don't include the trailing newline.
+ if self.selection.type == SelectionType.LINES and cut_text.endswith("\n"):
+ cut_text = cut_text[:-1]
+
+ return (
+ Document(text=remaining_text, cursor_position=new_cursor_position),
+ ClipboardData(cut_text, self.selection.type),
+ )
+ else:
+ return self, ClipboardData("")
+
+ def paste_clipboard_data(
+ self,
+ data: ClipboardData,
+ paste_mode: PasteMode = PasteMode.EMACS,
+ count: int = 1,
+ ) -> Document:
+ """
+ Return a new :class:`.Document` instance which contains the result if
+ we would paste this data at the current cursor position.
+
+ :param paste_mode: Where to paste. (Before/after/emacs.)
+ :param count: When >1, Paste multiple times.
+ """
+ before = paste_mode == PasteMode.VI_BEFORE
+ after = paste_mode == PasteMode.VI_AFTER
+
+ if data.type == SelectionType.CHARACTERS:
+ if after:
+ new_text = (
+ self.text[: self.cursor_position + 1]
+ + data.text * count
+ + self.text[self.cursor_position + 1 :]
+ )
+ else:
+ new_text = (
+ self.text_before_cursor + data.text * count + self.text_after_cursor
+ )
+
+ new_cursor_position = self.cursor_position + len(data.text) * count
+ if before:
+ new_cursor_position -= 1
+
+ elif data.type == SelectionType.LINES:
+ l = self.cursor_position_row
+ if before:
+ lines = self.lines[:l] + [data.text] * count + self.lines[l:]
+ new_text = "\n".join(lines)
+ new_cursor_position = len("".join(self.lines[:l])) + l
+ else:
+ lines = self.lines[: l + 1] + [data.text] * count + self.lines[l + 1 :]
+ new_cursor_position = len("".join(self.lines[: l + 1])) + l + 1
+ new_text = "\n".join(lines)
+
+ elif data.type == SelectionType.BLOCK:
+ lines = self.lines[:]
+ start_line = self.cursor_position_row
+ start_column = self.cursor_position_col + (0 if before else 1)
+
+ for i, line in enumerate(data.text.split("\n")):
+ index = i + start_line
+ if index >= len(lines):
+ lines.append("")
+
+ lines[index] = lines[index].ljust(start_column)
+ lines[index] = (
+ lines[index][:start_column]
+ + line * count
+ + lines[index][start_column:]
+ )
+
+ new_text = "\n".join(lines)
+ new_cursor_position = self.cursor_position + (0 if before else 1)
+
+ return Document(text=new_text, cursor_position=new_cursor_position)
+
+ def empty_line_count_at_the_end(self) -> int:
+ """
+ Return number of empty lines at the end of the document.
+ """
+ count = 0
+ for line in self.lines[::-1]:
+ if not line or line.isspace():
+ count += 1
+ else:
+ break
+
+ return count
+
+ def start_of_paragraph(self, count: int = 1, before: bool = False) -> int:
+ """
+ Return the start of the current paragraph. (Relative cursor position.)
+ """
+
+ def match_func(text: str) -> bool:
+ return not text or text.isspace()
+
+ line_index = self.find_previous_matching_line(
+ match_func=match_func, count=count
+ )
+
+ if line_index:
+ add = 0 if before else 1
+ return min(0, self.get_cursor_up_position(count=-line_index) + add)
+ else:
+ return -self.cursor_position
+
+ def end_of_paragraph(self, count: int = 1, after: bool = False) -> int:
+ """
+ Return the end of the current paragraph. (Relative cursor position.)
+ """
+
+ def match_func(text: str) -> bool:
+ return not text or text.isspace()
+
+ line_index = self.find_next_matching_line(match_func=match_func, count=count)
+
+ if line_index:
+ add = 0 if after else 1
+ return max(0, self.get_cursor_down_position(count=line_index) - add)
+ else:
+ return len(self.text_after_cursor)
+
+ # Modifiers.
+
+ def insert_after(self, text: str) -> Document:
+ """
+ Create a new document, with this text inserted after the buffer.
+ It keeps selection ranges and cursor position in sync.
+ """
+ return Document(
+ text=self.text + text,
+ cursor_position=self.cursor_position,
+ selection=self.selection,
+ )
+
+ def insert_before(self, text: str) -> Document:
+ """
+ Create a new document, with this text inserted before the buffer.
+ It keeps selection ranges and cursor position in sync.
+ """
+ selection_state = self.selection
+
+ if selection_state:
+ selection_state = SelectionState(
+ original_cursor_position=selection_state.original_cursor_position
+ + len(text),
+ type=selection_state.type,
+ )
+
+ return Document(
+ text=text + self.text,
+ cursor_position=self.cursor_position + len(text),
+ selection=selection_state,
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/enums.py b/lib/python3.12/site-packages/prompt_toolkit/enums.py
new file mode 100644
index 0000000000000000000000000000000000000000..da036337115df82cfbf027616c7ae8ecb7703dd9
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/enums.py
@@ -0,0 +1,19 @@
+from __future__ import annotations
+
+from enum import Enum
+
+
+class EditingMode(Enum):
+ # The set of key bindings that is active.
+ VI = "VI"
+ EMACS = "EMACS"
+
+
+#: Name of the search buffer.
+SEARCH_BUFFER = "SEARCH_BUFFER"
+
+#: Name of the default buffer.
+DEFAULT_BUFFER = "DEFAULT_BUFFER"
+
+#: Name of the system buffer.
+SYSTEM_BUFFER = "SYSTEM_BUFFER"
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/filters/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..556ed8846b84ad636e58e26621b24e98a103e685
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/filters/__init__.py
@@ -0,0 +1,71 @@
+"""
+Filters decide whether something is active or not (they decide about a boolean
+state). This is used to enable/disable features, like key bindings, parts of
+the layout and other stuff. For instance, we could have a `HasSearch` filter
+attached to some part of the layout, in order to show that part of the user
+interface only while the user is searching.
+
+Filters are made to avoid having to attach callbacks to all event in order to
+propagate state. However, they are lazy, they don't automatically propagate the
+state of what they are observing. Only when a filter is called (it's actually a
+callable), it will calculate its value. So, its not really reactive
+programming, but it's made to fit for this framework.
+
+Filters can be chained using ``&`` and ``|`` operations, and inverted using the
+``~`` operator, for instance::
+
+ filter = has_focus('default') & ~ has_selection
+"""
+
+from __future__ import annotations
+
+from .app import *
+from .base import Always, Condition, Filter, FilterOrBool, Never
+from .cli import *
+from .utils import is_true, to_filter
+
+__all__ = [
+ # app
+ "has_arg",
+ "has_completions",
+ "completion_is_selected",
+ "has_focus",
+ "buffer_has_focus",
+ "has_selection",
+ "has_validation_error",
+ "is_done",
+ "is_read_only",
+ "is_multiline",
+ "renderer_height_is_known",
+ "in_editing_mode",
+ "in_paste_mode",
+ "vi_mode",
+ "vi_navigation_mode",
+ "vi_insert_mode",
+ "vi_insert_multiple_mode",
+ "vi_replace_mode",
+ "vi_selection_mode",
+ "vi_waiting_for_text_object_mode",
+ "vi_digraph_mode",
+ "vi_recording_macro",
+ "emacs_mode",
+ "emacs_insert_mode",
+ "emacs_selection_mode",
+ "shift_selection_mode",
+ "is_searching",
+ "control_is_searchable",
+ "vi_search_direction_reversed",
+ # base.
+ "Filter",
+ "Never",
+ "Always",
+ "Condition",
+ "FilterOrBool",
+ # utils.
+ "is_true",
+ "to_filter",
+]
+
+from .cli import __all__ as cli_all
+
+__all__.extend(cli_all)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..624eddae950cd1985c6c59c6032e10c1a68e10dd
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/app.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/app.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..46580550d17b3ca7f36f4713b3277f2e32ee0a23
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/app.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/base.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/base.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..313958e8cd19e8f65cd8c199e031e749a6fa784c
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/base.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/cli.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/cli.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..90564806feecdeccca4ba414d1acd8d5ede63c1b
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/cli.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/utils.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/utils.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..102cce8aa55310360adf42bed441c7e2c2aa3595
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/filters/__pycache__/utils.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/app.py b/lib/python3.12/site-packages/prompt_toolkit/filters/app.py
new file mode 100644
index 0000000000000000000000000000000000000000..b9cc6111c90bdf040391fb591844b3ea1ecfef3e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/filters/app.py
@@ -0,0 +1,419 @@
+"""
+Filters that accept a `Application` as argument.
+"""
+
+from __future__ import annotations
+
+from typing import TYPE_CHECKING
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.cache import memoized
+from prompt_toolkit.enums import EditingMode
+
+from .base import Condition
+
+if TYPE_CHECKING:
+ from prompt_toolkit.layout.layout import FocusableElement
+
+
+__all__ = [
+ "has_arg",
+ "has_completions",
+ "completion_is_selected",
+ "has_focus",
+ "buffer_has_focus",
+ "has_selection",
+ "has_suggestion",
+ "has_validation_error",
+ "is_done",
+ "is_read_only",
+ "is_multiline",
+ "renderer_height_is_known",
+ "in_editing_mode",
+ "in_paste_mode",
+ "vi_mode",
+ "vi_navigation_mode",
+ "vi_insert_mode",
+ "vi_insert_multiple_mode",
+ "vi_replace_mode",
+ "vi_selection_mode",
+ "vi_waiting_for_text_object_mode",
+ "vi_digraph_mode",
+ "vi_recording_macro",
+ "emacs_mode",
+ "emacs_insert_mode",
+ "emacs_selection_mode",
+ "shift_selection_mode",
+ "is_searching",
+ "control_is_searchable",
+ "vi_search_direction_reversed",
+]
+
+
+# NOTE: `has_focus` below should *not* be `memoized`. It can reference any user
+# control. For instance, if we would continuously create new
+# `PromptSession` instances, then previous instances won't be released,
+# because this memoize (which caches results in the global scope) will
+# still refer to each instance.
+def has_focus(value: FocusableElement) -> Condition:
+ """
+ Enable when this buffer has the focus.
+ """
+ from prompt_toolkit.buffer import Buffer
+ from prompt_toolkit.layout import walk
+ from prompt_toolkit.layout.containers import Window, to_container
+ from prompt_toolkit.layout.controls import UIControl
+
+ if isinstance(value, str):
+
+ def test() -> bool:
+ return get_app().current_buffer.name == value
+
+ elif isinstance(value, Buffer):
+
+ def test() -> bool:
+ return get_app().current_buffer == value
+
+ elif isinstance(value, UIControl):
+
+ def test() -> bool:
+ return get_app().layout.current_control == value
+
+ else:
+ value = to_container(value)
+
+ if isinstance(value, Window):
+
+ def test() -> bool:
+ return get_app().layout.current_window == value
+
+ else:
+
+ def test() -> bool:
+ # Consider focused when any window inside this container is
+ # focused.
+ current_window = get_app().layout.current_window
+
+ for c in walk(value):
+ if isinstance(c, Window) and c == current_window:
+ return True
+ return False
+
+ @Condition
+ def has_focus_filter() -> bool:
+ return test()
+
+ return has_focus_filter
+
+
+@Condition
+def buffer_has_focus() -> bool:
+ """
+ Enabled when the currently focused control is a `BufferControl`.
+ """
+ return get_app().layout.buffer_has_focus
+
+
+@Condition
+def has_selection() -> bool:
+ """
+ Enable when the current buffer has a selection.
+ """
+ return bool(get_app().current_buffer.selection_state)
+
+
+@Condition
+def has_suggestion() -> bool:
+ """
+ Enable when the current buffer has a suggestion.
+ """
+ buffer = get_app().current_buffer
+ return buffer.suggestion is not None and buffer.suggestion.text != ""
+
+
+@Condition
+def has_completions() -> bool:
+ """
+ Enable when the current buffer has completions.
+ """
+ state = get_app().current_buffer.complete_state
+ return state is not None and len(state.completions) > 0
+
+
+@Condition
+def completion_is_selected() -> bool:
+ """
+ True when the user selected a completion.
+ """
+ complete_state = get_app().current_buffer.complete_state
+ return complete_state is not None and complete_state.current_completion is not None
+
+
+@Condition
+def is_read_only() -> bool:
+ """
+ True when the current buffer is read only.
+ """
+ return get_app().current_buffer.read_only()
+
+
+@Condition
+def is_multiline() -> bool:
+ """
+ True when the current buffer has been marked as multiline.
+ """
+ return get_app().current_buffer.multiline()
+
+
+@Condition
+def has_validation_error() -> bool:
+ "Current buffer has validation error."
+ return get_app().current_buffer.validation_error is not None
+
+
+@Condition
+def has_arg() -> bool:
+ "Enable when the input processor has an 'arg'."
+ return get_app().key_processor.arg is not None
+
+
+@Condition
+def is_done() -> bool:
+ """
+ True when the CLI is returning, aborting or exiting.
+ """
+ return get_app().is_done
+
+
+@Condition
+def renderer_height_is_known() -> bool:
+ """
+ Only True when the renderer knows it's real height.
+
+ (On VT100 terminals, we have to wait for a CPR response, before we can be
+ sure of the available height between the cursor position and the bottom of
+ the terminal. And usually it's nicer to wait with drawing bottom toolbars
+ until we receive the height, in order to avoid flickering -- first drawing
+ somewhere in the middle, and then again at the bottom.)
+ """
+ return get_app().renderer.height_is_known
+
+
+@memoized()
+def in_editing_mode(editing_mode: EditingMode) -> Condition:
+ """
+ Check whether a given editing mode is active. (Vi or Emacs.)
+ """
+
+ @Condition
+ def in_editing_mode_filter() -> bool:
+ return get_app().editing_mode == editing_mode
+
+ return in_editing_mode_filter
+
+
+@Condition
+def in_paste_mode() -> bool:
+ return get_app().paste_mode()
+
+
+@Condition
+def vi_mode() -> bool:
+ return get_app().editing_mode == EditingMode.VI
+
+
+@Condition
+def vi_navigation_mode() -> bool:
+ """
+ Active when the set for Vi navigation key bindings are active.
+ """
+ from prompt_toolkit.key_binding.vi_state import InputMode
+
+ app = get_app()
+
+ if (
+ app.editing_mode != EditingMode.VI
+ or app.vi_state.operator_func
+ or app.vi_state.waiting_for_digraph
+ or app.current_buffer.selection_state
+ ):
+ return False
+
+ return (
+ app.vi_state.input_mode == InputMode.NAVIGATION
+ or app.vi_state.temporary_navigation_mode
+ or app.current_buffer.read_only()
+ )
+
+
+@Condition
+def vi_insert_mode() -> bool:
+ from prompt_toolkit.key_binding.vi_state import InputMode
+
+ app = get_app()
+
+ if (
+ app.editing_mode != EditingMode.VI
+ or app.vi_state.operator_func
+ or app.vi_state.waiting_for_digraph
+ or app.current_buffer.selection_state
+ or app.vi_state.temporary_navigation_mode
+ or app.current_buffer.read_only()
+ ):
+ return False
+
+ return app.vi_state.input_mode == InputMode.INSERT
+
+
+@Condition
+def vi_insert_multiple_mode() -> bool:
+ from prompt_toolkit.key_binding.vi_state import InputMode
+
+ app = get_app()
+
+ if (
+ app.editing_mode != EditingMode.VI
+ or app.vi_state.operator_func
+ or app.vi_state.waiting_for_digraph
+ or app.current_buffer.selection_state
+ or app.vi_state.temporary_navigation_mode
+ or app.current_buffer.read_only()
+ ):
+ return False
+
+ return app.vi_state.input_mode == InputMode.INSERT_MULTIPLE
+
+
+@Condition
+def vi_replace_mode() -> bool:
+ from prompt_toolkit.key_binding.vi_state import InputMode
+
+ app = get_app()
+
+ if (
+ app.editing_mode != EditingMode.VI
+ or app.vi_state.operator_func
+ or app.vi_state.waiting_for_digraph
+ or app.current_buffer.selection_state
+ or app.vi_state.temporary_navigation_mode
+ or app.current_buffer.read_only()
+ ):
+ return False
+
+ return app.vi_state.input_mode == InputMode.REPLACE
+
+
+@Condition
+def vi_replace_single_mode() -> bool:
+ from prompt_toolkit.key_binding.vi_state import InputMode
+
+ app = get_app()
+
+ if (
+ app.editing_mode != EditingMode.VI
+ or app.vi_state.operator_func
+ or app.vi_state.waiting_for_digraph
+ or app.current_buffer.selection_state
+ or app.vi_state.temporary_navigation_mode
+ or app.current_buffer.read_only()
+ ):
+ return False
+
+ return app.vi_state.input_mode == InputMode.REPLACE_SINGLE
+
+
+@Condition
+def vi_selection_mode() -> bool:
+ app = get_app()
+ if app.editing_mode != EditingMode.VI:
+ return False
+
+ return bool(app.current_buffer.selection_state)
+
+
+@Condition
+def vi_waiting_for_text_object_mode() -> bool:
+ app = get_app()
+ if app.editing_mode != EditingMode.VI:
+ return False
+
+ return app.vi_state.operator_func is not None
+
+
+@Condition
+def vi_digraph_mode() -> bool:
+ app = get_app()
+ if app.editing_mode != EditingMode.VI:
+ return False
+
+ return app.vi_state.waiting_for_digraph
+
+
+@Condition
+def vi_recording_macro() -> bool:
+ "When recording a Vi macro."
+ app = get_app()
+ if app.editing_mode != EditingMode.VI:
+ return False
+
+ return app.vi_state.recording_register is not None
+
+
+@Condition
+def emacs_mode() -> bool:
+ "When the Emacs bindings are active."
+ return get_app().editing_mode == EditingMode.EMACS
+
+
+@Condition
+def emacs_insert_mode() -> bool:
+ app = get_app()
+ if (
+ app.editing_mode != EditingMode.EMACS
+ or app.current_buffer.selection_state
+ or app.current_buffer.read_only()
+ ):
+ return False
+ return True
+
+
+@Condition
+def emacs_selection_mode() -> bool:
+ app = get_app()
+ return bool(
+ app.editing_mode == EditingMode.EMACS and app.current_buffer.selection_state
+ )
+
+
+@Condition
+def shift_selection_mode() -> bool:
+ app = get_app()
+ return bool(
+ app.current_buffer.selection_state
+ and app.current_buffer.selection_state.shift_mode
+ )
+
+
+@Condition
+def is_searching() -> bool:
+ "When we are searching."
+ app = get_app()
+ return app.layout.is_searching
+
+
+@Condition
+def control_is_searchable() -> bool:
+ "When the current UIControl is searchable."
+ from prompt_toolkit.layout.controls import BufferControl
+
+ control = get_app().layout.current_control
+
+ return (
+ isinstance(control, BufferControl) and control.search_buffer_control is not None
+ )
+
+
+@Condition
+def vi_search_direction_reversed() -> bool:
+ "When the '/' and '?' key bindings for Vi-style searching have been reversed."
+ return get_app().reverse_vi_search_direction()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/base.py b/lib/python3.12/site-packages/prompt_toolkit/filters/base.py
new file mode 100644
index 0000000000000000000000000000000000000000..cd95424dc3862e14cfb57f079d7dfdcbb255bd86
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/filters/base.py
@@ -0,0 +1,260 @@
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from typing import Callable, Iterable, Union
+
+__all__ = ["Filter", "Never", "Always", "Condition", "FilterOrBool"]
+
+
+class Filter(metaclass=ABCMeta):
+ """
+ Base class for any filter to activate/deactivate a feature, depending on a
+ condition.
+
+ The return value of ``__call__`` will tell if the feature should be active.
+ """
+
+ def __init__(self) -> None:
+ self._and_cache: dict[Filter, Filter] = {}
+ self._or_cache: dict[Filter, Filter] = {}
+ self._invert_result: Filter | None = None
+
+ @abstractmethod
+ def __call__(self) -> bool:
+ """
+ The actual call to evaluate the filter.
+ """
+ return True
+
+ def __and__(self, other: Filter) -> Filter:
+ """
+ Chaining of filters using the & operator.
+ """
+ assert isinstance(other, Filter), f"Expecting filter, got {other!r}"
+
+ if isinstance(other, Always):
+ return self
+ if isinstance(other, Never):
+ return other
+
+ if other in self._and_cache:
+ return self._and_cache[other]
+
+ result = _AndList.create([self, other])
+ self._and_cache[other] = result
+ return result
+
+ def __or__(self, other: Filter) -> Filter:
+ """
+ Chaining of filters using the | operator.
+ """
+ assert isinstance(other, Filter), f"Expecting filter, got {other!r}"
+
+ if isinstance(other, Always):
+ return other
+ if isinstance(other, Never):
+ return self
+
+ if other in self._or_cache:
+ return self._or_cache[other]
+
+ result = _OrList.create([self, other])
+ self._or_cache[other] = result
+ return result
+
+ def __invert__(self) -> Filter:
+ """
+ Inverting of filters using the ~ operator.
+ """
+ if self._invert_result is None:
+ self._invert_result = _Invert(self)
+
+ return self._invert_result
+
+ def __bool__(self) -> None:
+ """
+ By purpose, we don't allow bool(...) operations directly on a filter,
+ because the meaning is ambiguous.
+
+ Executing a filter has to be done always by calling it. Providing
+ defaults for `None` values should be done through an `is None` check
+ instead of for instance ``filter1 or Always()``.
+ """
+ raise ValueError(
+ "The truth value of a Filter is ambiguous. Instead, call it as a function."
+ )
+
+
+def _remove_duplicates(filters: list[Filter]) -> list[Filter]:
+ result = []
+ for f in filters:
+ if f not in result:
+ result.append(f)
+ return result
+
+
+class _AndList(Filter):
+ """
+ Result of &-operation between several filters.
+ """
+
+ def __init__(self, filters: list[Filter]) -> None:
+ super().__init__()
+ self.filters = filters
+
+ @classmethod
+ def create(cls, filters: Iterable[Filter]) -> Filter:
+ """
+ Create a new filter by applying an `&` operator between them.
+
+ If there's only one unique filter in the given iterable, it will return
+ that one filter instead of an `_AndList`.
+ """
+ filters_2: list[Filter] = []
+
+ for f in filters:
+ if isinstance(f, _AndList): # Turn nested _AndLists into one.
+ filters_2.extend(f.filters)
+ else:
+ filters_2.append(f)
+
+ # Remove duplicates. This could speed up execution, and doesn't make a
+ # difference for the evaluation.
+ filters = _remove_duplicates(filters_2)
+
+ # If only one filter is left, return that without wrapping into an
+ # `_AndList`.
+ if len(filters) == 1:
+ return filters[0]
+
+ return cls(filters)
+
+ def __call__(self) -> bool:
+ return all(f() for f in self.filters)
+
+ def __repr__(self) -> str:
+ return "&".join(repr(f) for f in self.filters)
+
+
+class _OrList(Filter):
+ """
+ Result of |-operation between several filters.
+ """
+
+ def __init__(self, filters: list[Filter]) -> None:
+ super().__init__()
+ self.filters = filters
+
+ @classmethod
+ def create(cls, filters: Iterable[Filter]) -> Filter:
+ """
+ Create a new filter by applying an `|` operator between them.
+
+ If there's only one unique filter in the given iterable, it will return
+ that one filter instead of an `_OrList`.
+ """
+ filters_2: list[Filter] = []
+
+ for f in filters:
+ if isinstance(f, _OrList): # Turn nested _AndLists into one.
+ filters_2.extend(f.filters)
+ else:
+ filters_2.append(f)
+
+ # Remove duplicates. This could speed up execution, and doesn't make a
+ # difference for the evaluation.
+ filters = _remove_duplicates(filters_2)
+
+ # If only one filter is left, return that without wrapping into an
+ # `_AndList`.
+ if len(filters) == 1:
+ return filters[0]
+
+ return cls(filters)
+
+ def __call__(self) -> bool:
+ return any(f() for f in self.filters)
+
+ def __repr__(self) -> str:
+ return "|".join(repr(f) for f in self.filters)
+
+
+class _Invert(Filter):
+ """
+ Negation of another filter.
+ """
+
+ def __init__(self, filter: Filter) -> None:
+ super().__init__()
+ self.filter = filter
+
+ def __call__(self) -> bool:
+ return not self.filter()
+
+ def __repr__(self) -> str:
+ return f"~{self.filter!r}"
+
+
+class Always(Filter):
+ """
+ Always enable feature.
+ """
+
+ def __call__(self) -> bool:
+ return True
+
+ def __or__(self, other: Filter) -> Filter:
+ return self
+
+ def __and__(self, other: Filter) -> Filter:
+ return other
+
+ def __invert__(self) -> Never:
+ return Never()
+
+
+class Never(Filter):
+ """
+ Never enable feature.
+ """
+
+ def __call__(self) -> bool:
+ return False
+
+ def __and__(self, other: Filter) -> Filter:
+ return self
+
+ def __or__(self, other: Filter) -> Filter:
+ return other
+
+ def __invert__(self) -> Always:
+ return Always()
+
+
+class Condition(Filter):
+ """
+ Turn any callable into a Filter. The callable is supposed to not take any
+ arguments.
+
+ This can be used as a decorator::
+
+ @Condition
+ def feature_is_active(): # `feature_is_active` becomes a Filter.
+ return True
+
+ :param func: Callable which takes no inputs and returns a boolean.
+ """
+
+ def __init__(self, func: Callable[[], bool]) -> None:
+ super().__init__()
+ self.func = func
+
+ def __call__(self) -> bool:
+ return self.func()
+
+ def __repr__(self) -> str:
+ return f"Condition({self.func!r})"
+
+
+# Often used as type annotation.
+FilterOrBool = Union[Filter, bool]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/cli.py b/lib/python3.12/site-packages/prompt_toolkit/filters/cli.py
new file mode 100644
index 0000000000000000000000000000000000000000..902fbaae981296d8bd8143db58463b3a3f9ddd10
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/filters/cli.py
@@ -0,0 +1,65 @@
+"""
+For backwards-compatibility. keep this file.
+(Many people are going to have key bindings that rely on this file.)
+"""
+
+from __future__ import annotations
+
+from .app import *
+
+__all__ = [
+ # Old names.
+ "HasArg",
+ "HasCompletions",
+ "HasFocus",
+ "HasSelection",
+ "HasValidationError",
+ "IsDone",
+ "IsReadOnly",
+ "IsMultiline",
+ "RendererHeightIsKnown",
+ "InEditingMode",
+ "InPasteMode",
+ "ViMode",
+ "ViNavigationMode",
+ "ViInsertMode",
+ "ViInsertMultipleMode",
+ "ViReplaceMode",
+ "ViSelectionMode",
+ "ViWaitingForTextObjectMode",
+ "ViDigraphMode",
+ "EmacsMode",
+ "EmacsInsertMode",
+ "EmacsSelectionMode",
+ "IsSearching",
+ "HasSearch",
+ "ControlIsSearchable",
+]
+
+# Keep the original classnames for backwards compatibility.
+HasValidationError = lambda: has_validation_error
+HasArg = lambda: has_arg
+IsDone = lambda: is_done
+RendererHeightIsKnown = lambda: renderer_height_is_known
+ViNavigationMode = lambda: vi_navigation_mode
+InPasteMode = lambda: in_paste_mode
+EmacsMode = lambda: emacs_mode
+EmacsInsertMode = lambda: emacs_insert_mode
+ViMode = lambda: vi_mode
+IsSearching = lambda: is_searching
+HasSearch = lambda: is_searching
+ControlIsSearchable = lambda: control_is_searchable
+EmacsSelectionMode = lambda: emacs_selection_mode
+ViDigraphMode = lambda: vi_digraph_mode
+ViWaitingForTextObjectMode = lambda: vi_waiting_for_text_object_mode
+ViSelectionMode = lambda: vi_selection_mode
+ViReplaceMode = lambda: vi_replace_mode
+ViInsertMultipleMode = lambda: vi_insert_multiple_mode
+ViInsertMode = lambda: vi_insert_mode
+HasSelection = lambda: has_selection
+HasCompletions = lambda: has_completions
+IsReadOnly = lambda: is_read_only
+IsMultiline = lambda: is_multiline
+
+HasFocus = has_focus # No lambda here! (Has_focus is callable that returns a callable.)
+InEditingMode = in_editing_mode
diff --git a/lib/python3.12/site-packages/prompt_toolkit/filters/utils.py b/lib/python3.12/site-packages/prompt_toolkit/filters/utils.py
new file mode 100644
index 0000000000000000000000000000000000000000..20e00ee09e3becc6d1e39621e63c1fac86ebe130
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/filters/utils.py
@@ -0,0 +1,41 @@
+from __future__ import annotations
+
+from .base import Always, Filter, FilterOrBool, Never
+
+__all__ = [
+ "to_filter",
+ "is_true",
+]
+
+
+_always = Always()
+_never = Never()
+
+
+_bool_to_filter: dict[bool, Filter] = {
+ True: _always,
+ False: _never,
+}
+
+
+def to_filter(bool_or_filter: FilterOrBool) -> Filter:
+ """
+ Accept both booleans and Filters as input and
+ turn it into a Filter.
+ """
+ if isinstance(bool_or_filter, bool):
+ return _bool_to_filter[bool_or_filter]
+
+ if isinstance(bool_or_filter, Filter):
+ return bool_or_filter
+
+ raise TypeError(f"Expecting a bool or a Filter instance. Got {bool_or_filter!r}")
+
+
+def is_true(value: FilterOrBool) -> bool:
+ """
+ Test whether `value` is True. In case of a Filter, call it.
+
+ :param value: Boolean or `Filter` instance.
+ """
+ return to_filter(value)()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/history.py b/lib/python3.12/site-packages/prompt_toolkit/history.py
new file mode 100644
index 0000000000000000000000000000000000000000..2d497a0209c4047ef2170ab992d8c299f7267822
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/history.py
@@ -0,0 +1,306 @@
+"""
+Implementations for the history of a `Buffer`.
+
+NOTE: There is no `DynamicHistory`:
+ This doesn't work well, because the `Buffer` needs to be able to attach
+ an event handler to the event when a history entry is loaded. This
+ loading can be done asynchronously and making the history swappable would
+ probably break this.
+"""
+
+from __future__ import annotations
+
+import datetime
+import os
+import threading
+from abc import ABCMeta, abstractmethod
+from asyncio import get_running_loop
+from typing import AsyncGenerator, Iterable, Sequence, Union
+
+__all__ = [
+ "History",
+ "ThreadedHistory",
+ "DummyHistory",
+ "FileHistory",
+ "InMemoryHistory",
+]
+
+
+class History(metaclass=ABCMeta):
+ """
+ Base ``History`` class.
+
+ This also includes abstract methods for loading/storing history.
+ """
+
+ def __init__(self) -> None:
+ # In memory storage for strings.
+ self._loaded = False
+
+ # History that's loaded already, in reverse order. Latest, most recent
+ # item first.
+ self._loaded_strings: list[str] = []
+
+ #
+ # Methods expected by `Buffer`.
+ #
+
+ async def load(self) -> AsyncGenerator[str, None]:
+ """
+ Load the history and yield all the entries in reverse order (latest,
+ most recent history entry first).
+
+ This method can be called multiple times from the `Buffer` to
+ repopulate the history when prompting for a new input. So we are
+ responsible here for both caching, and making sure that strings that
+ were were appended to the history will be incorporated next time this
+ method is called.
+ """
+ if not self._loaded:
+ self._loaded_strings = list(self.load_history_strings())
+ self._loaded = True
+
+ for item in self._loaded_strings:
+ yield item
+
+ def get_strings(self) -> list[str]:
+ """
+ Get the strings from the history that are loaded so far.
+ (In order. Oldest item first.)
+ """
+ return self._loaded_strings[::-1]
+
+ def append_string(self, string: str) -> None:
+ "Add string to the history."
+ self._loaded_strings.insert(0, string)
+ self.store_string(string)
+
+ #
+ # Implementation for specific backends.
+ #
+
+ @abstractmethod
+ def load_history_strings(self) -> Iterable[str]:
+ """
+ This should be a generator that yields `str` instances.
+
+ It should yield the most recent items first, because they are the most
+ important. (The history can already be used, even when it's only
+ partially loaded.)
+ """
+ while False:
+ yield
+
+ @abstractmethod
+ def store_string(self, string: str) -> None:
+ """
+ Store the string in persistent storage.
+ """
+
+
+class ThreadedHistory(History):
+ """
+ Wrapper around `History` implementations that run the `load()` generator in
+ a thread.
+
+ Use this to increase the start-up time of prompt_toolkit applications.
+ History entries are available as soon as they are loaded. We don't have to
+ wait for everything to be loaded.
+ """
+
+ def __init__(self, history: History) -> None:
+ super().__init__()
+
+ self.history = history
+
+ self._load_thread: threading.Thread | None = None
+
+ # Lock for accessing/manipulating `_loaded_strings` and `_loaded`
+ # together in a consistent state.
+ self._lock = threading.Lock()
+
+ # Events created by each `load()` call. Used to wait for new history
+ # entries from the loader thread.
+ self._string_load_events: list[threading.Event] = []
+
+ async def load(self) -> AsyncGenerator[str, None]:
+ """
+ Like `History.load(), but call `self.load_history_strings()` in a
+ background thread.
+ """
+ # Start the load thread, if this is called for the first time.
+ if not self._load_thread:
+ self._load_thread = threading.Thread(
+ target=self._in_load_thread,
+ daemon=True,
+ )
+ self._load_thread.start()
+
+ # Consume the `_loaded_strings` list, using asyncio.
+ loop = get_running_loop()
+
+ # Create threading Event so that we can wait for new items.
+ event = threading.Event()
+ event.set()
+ self._string_load_events.append(event)
+
+ items_yielded = 0
+
+ try:
+ while True:
+ # Wait for new items to be available.
+ # (Use a timeout, because the executor thread is not a daemon
+ # thread. The "slow-history.py" example would otherwise hang if
+ # Control-C is pressed before the history is fully loaded,
+ # because there's still this non-daemon executor thread waiting
+ # for this event.)
+ got_timeout = await loop.run_in_executor(
+ None, lambda: event.wait(timeout=0.5)
+ )
+ if not got_timeout:
+ continue
+
+ # Read new items (in lock).
+ def in_executor() -> tuple[list[str], bool]:
+ with self._lock:
+ new_items = self._loaded_strings[items_yielded:]
+ done = self._loaded
+ event.clear()
+ return new_items, done
+
+ new_items, done = await loop.run_in_executor(None, in_executor)
+
+ items_yielded += len(new_items)
+
+ for item in new_items:
+ yield item
+
+ if done:
+ break
+ finally:
+ self._string_load_events.remove(event)
+
+ def _in_load_thread(self) -> None:
+ try:
+ # Start with an empty list. In case `append_string()` was called
+ # before `load()` happened. Then `.store_string()` will have
+ # written these entries back to disk and we will reload it.
+ self._loaded_strings = []
+
+ for item in self.history.load_history_strings():
+ with self._lock:
+ self._loaded_strings.append(item)
+
+ for event in self._string_load_events:
+ event.set()
+ finally:
+ with self._lock:
+ self._loaded = True
+ for event in self._string_load_events:
+ event.set()
+
+ def append_string(self, string: str) -> None:
+ with self._lock:
+ self._loaded_strings.insert(0, string)
+ self.store_string(string)
+
+ # All of the following are proxied to `self.history`.
+
+ def load_history_strings(self) -> Iterable[str]:
+ return self.history.load_history_strings()
+
+ def store_string(self, string: str) -> None:
+ self.history.store_string(string)
+
+ def __repr__(self) -> str:
+ return f"ThreadedHistory({self.history!r})"
+
+
+class InMemoryHistory(History):
+ """
+ :class:`.History` class that keeps a list of all strings in memory.
+
+ In order to prepopulate the history, it's possible to call either
+ `append_string` for all items or pass a list of strings to `__init__` here.
+ """
+
+ def __init__(self, history_strings: Sequence[str] | None = None) -> None:
+ super().__init__()
+ # Emulating disk storage.
+ if history_strings is None:
+ self._storage = []
+ else:
+ self._storage = list(history_strings)
+
+ def load_history_strings(self) -> Iterable[str]:
+ yield from self._storage[::-1]
+
+ def store_string(self, string: str) -> None:
+ self._storage.append(string)
+
+
+class DummyHistory(History):
+ """
+ :class:`.History` object that doesn't remember anything.
+ """
+
+ def load_history_strings(self) -> Iterable[str]:
+ return []
+
+ def store_string(self, string: str) -> None:
+ pass
+
+ def append_string(self, string: str) -> None:
+ # Don't remember this.
+ pass
+
+
+_StrOrBytesPath = Union[str, bytes, "os.PathLike[str]", "os.PathLike[bytes]"]
+
+
+class FileHistory(History):
+ """
+ :class:`.History` class that stores all strings in a file.
+ """
+
+ def __init__(self, filename: _StrOrBytesPath) -> None:
+ self.filename = filename
+ super().__init__()
+
+ def load_history_strings(self) -> Iterable[str]:
+ strings: list[str] = []
+ lines: list[str] = []
+
+ def add() -> None:
+ if lines:
+ # Join and drop trailing newline.
+ string = "".join(lines)[:-1]
+
+ strings.append(string)
+
+ if os.path.exists(self.filename):
+ with open(self.filename, "rb") as f:
+ for line_bytes in f:
+ line = line_bytes.decode("utf-8", errors="replace")
+
+ if line.startswith("+"):
+ lines.append(line[1:])
+ else:
+ add()
+ lines = []
+
+ add()
+
+ # Reverse the order, because newest items have to go first.
+ return reversed(strings)
+
+ def store_string(self, string: str) -> None:
+ # Save to file.
+ with open(self.filename, "ab") as f:
+
+ def write(t: str) -> None:
+ f.write(t.encode("utf-8"))
+
+ write(f"\n# {datetime.datetime.now()}\n")
+ for line in string.split("\n"):
+ write(f"+{line}\n")
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..c31746aba14fd18b4d2a6df8c474bcebddfdb057
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__init__.py
@@ -0,0 +1,22 @@
+from __future__ import annotations
+
+from .key_bindings import (
+ ConditionalKeyBindings,
+ DynamicKeyBindings,
+ KeyBindings,
+ KeyBindingsBase,
+ merge_key_bindings,
+)
+from .key_processor import KeyPress, KeyPressEvent
+
+__all__ = [
+ # key_bindings.
+ "ConditionalKeyBindings",
+ "DynamicKeyBindings",
+ "KeyBindings",
+ "KeyBindingsBase",
+ "merge_key_bindings",
+ # key_processor
+ "KeyPress",
+ "KeyPressEvent",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..040b972eef2f6bf4c467da91550c5206c73effc9
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/defaults.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/defaults.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..ed9232680e31e34781e27f5b4fdefd6036da52a0
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/defaults.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/digraphs.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/digraphs.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..4e668bc526040f248aa0197e63df44204e0b036a
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/digraphs.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/emacs_state.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/emacs_state.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..5feb462010905d18469973b89145fc6b1d8f0d38
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/emacs_state.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/key_bindings.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/key_bindings.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..af886942c73e34a29453235628b0e8116c0c145e
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/key_bindings.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/key_processor.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/key_processor.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..4c2e4c4bbb39541615afc32dbe146c1dc3d1ce7b
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/key_processor.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/vi_state.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/vi_state.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..921cff96804fb777722c1fb464dead4613024125
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/__pycache__/vi_state.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..053c7e76f87f9d185066aa5c809f9c3c44c0a262
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/auto_suggest.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/auto_suggest.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..01febbd068b8c6cad7ba73f86ea971a9f0250d00
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/auto_suggest.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/basic.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/basic.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..342c920b826e9f157426ef3afa45c53d197c27ea
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/basic.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/completion.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/completion.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..2021d9616439e900a4746c067da2b6634caf9b71
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/completion.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/cpr.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/cpr.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..ea3facec2d12decf102cf5aefcc6b863829f60dd
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/cpr.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/emacs.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/emacs.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..3ba3c5db5db1be52784a5e5f1d92379c40994cfc
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/emacs.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/focus.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/focus.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..587ff59d6829984ae3bdafcf5fb6b02a162e3b0c
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/focus.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/mouse.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/mouse.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..2544c2c57e0afaea829f31351fde65ca2d80dcd1
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/mouse.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/named_commands.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/named_commands.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..01f4986a8bdb5643a29288c198f24ad0f9043eba
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/named_commands.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/open_in_editor.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/open_in_editor.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..c6cb4e7293b9fdab1f713d8eea846a6e754a36ef
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/open_in_editor.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/page_navigation.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/page_navigation.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..8d817880b4bb9ecfd5476279dd0a12c9b0601d80
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/page_navigation.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/scroll.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/scroll.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..a401680308214efbac840ebe9cf14d02333d8907
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/scroll.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/search.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/search.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..46516cc0e0e03005a069b16fe1fba57000e42091
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/__pycache__/search.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/auto_suggest.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/auto_suggest.py
new file mode 100644
index 0000000000000000000000000000000000000000..b487f14ffe52ce020f2f2585c2df5629b6e3557b
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/auto_suggest.py
@@ -0,0 +1,66 @@
+"""
+Key bindings for auto suggestion (for fish-style auto suggestion).
+"""
+
+from __future__ import annotations
+
+import re
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.filters import Condition, emacs_mode
+from prompt_toolkit.key_binding.key_bindings import KeyBindings
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+
+__all__ = [
+ "load_auto_suggest_bindings",
+]
+
+E = KeyPressEvent
+
+
+def load_auto_suggest_bindings() -> KeyBindings:
+ """
+ Key bindings for accepting auto suggestion text.
+
+ (This has to come after the Vi bindings, because they also have an
+ implementation for the "right arrow", but we really want the suggestion
+ binding when a suggestion is available.)
+ """
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+
+ @Condition
+ def suggestion_available() -> bool:
+ app = get_app()
+ return (
+ app.current_buffer.suggestion is not None
+ and len(app.current_buffer.suggestion.text) > 0
+ and app.current_buffer.document.is_cursor_at_the_end
+ )
+
+ @handle("c-f", filter=suggestion_available)
+ @handle("c-e", filter=suggestion_available)
+ @handle("right", filter=suggestion_available)
+ def _accept(event: E) -> None:
+ """
+ Accept suggestion.
+ """
+ b = event.current_buffer
+ suggestion = b.suggestion
+
+ if suggestion:
+ b.insert_text(suggestion.text)
+
+ @handle("escape", "f", filter=suggestion_available & emacs_mode)
+ def _fill(event: E) -> None:
+ """
+ Fill partial suggestion.
+ """
+ b = event.current_buffer
+ suggestion = b.suggestion
+
+ if suggestion:
+ t = re.split(r"([^\s/]+(?:\s+|/))", suggestion.text)
+ b.insert_text(next(x for x in t if x))
+
+ return key_bindings
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/basic.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/basic.py
new file mode 100644
index 0000000000000000000000000000000000000000..ad18df981c4889585e6410ad950c063ffc54516f
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/basic.py
@@ -0,0 +1,257 @@
+# pylint: disable=function-redefined
+from __future__ import annotations
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.filters import (
+ Condition,
+ emacs_insert_mode,
+ has_selection,
+ in_paste_mode,
+ is_multiline,
+ vi_insert_mode,
+)
+from prompt_toolkit.key_binding.key_processor import KeyPress, KeyPressEvent
+from prompt_toolkit.keys import Keys
+
+from ..key_bindings import KeyBindings
+from .named_commands import get_by_name
+
+__all__ = [
+ "load_basic_bindings",
+]
+
+E = KeyPressEvent
+
+
+def if_no_repeat(event: E) -> bool:
+ """Callable that returns True when the previous event was delivered to
+ another handler."""
+ return not event.is_repeat
+
+
+@Condition
+def has_text_before_cursor() -> bool:
+ return bool(get_app().current_buffer.text)
+
+
+@Condition
+def in_quoted_insert() -> bool:
+ return get_app().quoted_insert
+
+
+def load_basic_bindings() -> KeyBindings:
+ key_bindings = KeyBindings()
+ insert_mode = vi_insert_mode | emacs_insert_mode
+ handle = key_bindings.add
+
+ @handle("c-a")
+ @handle("c-b")
+ @handle("c-c")
+ @handle("c-d")
+ @handle("c-e")
+ @handle("c-f")
+ @handle("c-g")
+ @handle("c-h")
+ @handle("c-i")
+ @handle("c-j")
+ @handle("c-k")
+ @handle("c-l")
+ @handle("c-m")
+ @handle("c-n")
+ @handle("c-o")
+ @handle("c-p")
+ @handle("c-q")
+ @handle("c-r")
+ @handle("c-s")
+ @handle("c-t")
+ @handle("c-u")
+ @handle("c-v")
+ @handle("c-w")
+ @handle("c-x")
+ @handle("c-y")
+ @handle("c-z")
+ @handle("f1")
+ @handle("f2")
+ @handle("f3")
+ @handle("f4")
+ @handle("f5")
+ @handle("f6")
+ @handle("f7")
+ @handle("f8")
+ @handle("f9")
+ @handle("f10")
+ @handle("f11")
+ @handle("f12")
+ @handle("f13")
+ @handle("f14")
+ @handle("f15")
+ @handle("f16")
+ @handle("f17")
+ @handle("f18")
+ @handle("f19")
+ @handle("f20")
+ @handle("f21")
+ @handle("f22")
+ @handle("f23")
+ @handle("f24")
+ @handle("c-@") # Also c-space.
+ @handle("c-\\")
+ @handle("c-]")
+ @handle("c-^")
+ @handle("c-_")
+ @handle("backspace")
+ @handle("up")
+ @handle("down")
+ @handle("right")
+ @handle("left")
+ @handle("s-up")
+ @handle("s-down")
+ @handle("s-right")
+ @handle("s-left")
+ @handle("home")
+ @handle("end")
+ @handle("s-home")
+ @handle("s-end")
+ @handle("delete")
+ @handle("s-delete")
+ @handle("c-delete")
+ @handle("pageup")
+ @handle("pagedown")
+ @handle("s-tab")
+ @handle("tab")
+ @handle("c-s-left")
+ @handle("c-s-right")
+ @handle("c-s-home")
+ @handle("c-s-end")
+ @handle("c-left")
+ @handle("c-right")
+ @handle("c-up")
+ @handle("c-down")
+ @handle("c-home")
+ @handle("c-end")
+ @handle("insert")
+ @handle("s-insert")
+ @handle("c-insert")
+ @handle("")
+ @handle(Keys.Ignore)
+ def _ignore(event: E) -> None:
+ """
+ First, for any of these keys, Don't do anything by default. Also don't
+ catch them in the 'Any' handler which will insert them as data.
+
+ If people want to insert these characters as a literal, they can always
+ do by doing a quoted insert. (ControlQ in emacs mode, ControlV in Vi
+ mode.)
+ """
+ pass
+
+ # Readline-style bindings.
+ handle("home")(get_by_name("beginning-of-line"))
+ handle("end")(get_by_name("end-of-line"))
+ handle("left")(get_by_name("backward-char"))
+ handle("right")(get_by_name("forward-char"))
+ handle("c-up")(get_by_name("previous-history"))
+ handle("c-down")(get_by_name("next-history"))
+ handle("c-l")(get_by_name("clear-screen"))
+
+ handle("c-k", filter=insert_mode)(get_by_name("kill-line"))
+ handle("c-u", filter=insert_mode)(get_by_name("unix-line-discard"))
+ handle("backspace", filter=insert_mode, save_before=if_no_repeat)(
+ get_by_name("backward-delete-char")
+ )
+ handle("delete", filter=insert_mode, save_before=if_no_repeat)(
+ get_by_name("delete-char")
+ )
+ handle("c-delete", filter=insert_mode, save_before=if_no_repeat)(
+ get_by_name("delete-char")
+ )
+ handle(Keys.Any, filter=insert_mode, save_before=if_no_repeat)(
+ get_by_name("self-insert")
+ )
+ handle("c-t", filter=insert_mode)(get_by_name("transpose-chars"))
+ handle("c-i", filter=insert_mode)(get_by_name("menu-complete"))
+ handle("s-tab", filter=insert_mode)(get_by_name("menu-complete-backward"))
+
+ # Control-W should delete, using whitespace as separator, while M-Del
+ # should delete using [^a-zA-Z0-9] as a boundary.
+ handle("c-w", filter=insert_mode)(get_by_name("unix-word-rubout"))
+
+ handle("pageup", filter=~has_selection)(get_by_name("previous-history"))
+ handle("pagedown", filter=~has_selection)(get_by_name("next-history"))
+
+ # CTRL keys.
+
+ handle("c-d", filter=has_text_before_cursor & insert_mode)(
+ get_by_name("delete-char")
+ )
+
+ @handle("enter", filter=insert_mode & is_multiline)
+ def _newline(event: E) -> None:
+ """
+ Newline (in case of multiline input.
+ """
+ event.current_buffer.newline(copy_margin=not in_paste_mode())
+
+ @handle("c-j")
+ def _newline2(event: E) -> None:
+ r"""
+ By default, handle \n as if it were a \r (enter).
+ (It appears that some terminals send \n instead of \r when pressing
+ enter. - at least the Linux subsystem for Windows.)
+ """
+ event.key_processor.feed(KeyPress(Keys.ControlM, "\r"), first=True)
+
+ # Delete the word before the cursor.
+
+ @handle("up")
+ def _go_up(event: E) -> None:
+ event.current_buffer.auto_up(count=event.arg)
+
+ @handle("down")
+ def _go_down(event: E) -> None:
+ event.current_buffer.auto_down(count=event.arg)
+
+ @handle("delete", filter=has_selection)
+ def _cut(event: E) -> None:
+ data = event.current_buffer.cut_selection()
+ event.app.clipboard.set_data(data)
+
+ # Global bindings.
+
+ @handle("c-z")
+ def _insert_ctrl_z(event: E) -> None:
+ """
+ By default, control-Z should literally insert Ctrl-Z.
+ (Ansi Ctrl-Z, code 26 in MSDOS means End-Of-File.
+ In a Python REPL for instance, it's possible to type
+ Control-Z followed by enter to quit.)
+
+ When the system bindings are loaded and suspend-to-background is
+ supported, that will override this binding.
+ """
+ event.current_buffer.insert_text(event.data)
+
+ @handle(Keys.BracketedPaste)
+ def _paste(event: E) -> None:
+ """
+ Pasting from clipboard.
+ """
+ data = event.data
+
+ # Be sure to use \n as line ending.
+ # Some terminals (Like iTerm2) seem to paste \r\n line endings in a
+ # bracketed paste. See: https://github.com/ipython/ipython/issues/9737
+ data = data.replace("\r\n", "\n")
+ data = data.replace("\r", "\n")
+
+ event.current_buffer.insert_text(data)
+
+ @handle(Keys.Any, filter=in_quoted_insert, eager=True)
+ def _insert_text(event: E) -> None:
+ """
+ Handle quoted insert.
+ """
+ event.current_buffer.insert_text(event.data, overwrite=False)
+ event.app.quoted_insert = False
+
+ return key_bindings
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/completion.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/completion.py
new file mode 100644
index 0000000000000000000000000000000000000000..8c5e005da8d14360c1575a792b1e1126f820bf1c
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/completion.py
@@ -0,0 +1,206 @@
+"""
+Key binding handlers for displaying completions.
+"""
+
+from __future__ import annotations
+
+import asyncio
+import math
+from typing import TYPE_CHECKING
+
+from prompt_toolkit.application.run_in_terminal import in_terminal
+from prompt_toolkit.completion import (
+ CompleteEvent,
+ Completion,
+ get_common_complete_suffix,
+)
+from prompt_toolkit.formatted_text import StyleAndTextTuples
+from prompt_toolkit.key_binding.key_bindings import KeyBindings
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+from prompt_toolkit.keys import Keys
+from prompt_toolkit.utils import get_cwidth
+
+if TYPE_CHECKING:
+ from prompt_toolkit.application import Application
+ from prompt_toolkit.shortcuts import PromptSession
+
+__all__ = [
+ "generate_completions",
+ "display_completions_like_readline",
+]
+
+E = KeyPressEvent
+
+
+def generate_completions(event: E) -> None:
+ r"""
+ Tab-completion: where the first tab completes the common suffix and the
+ second tab lists all the completions.
+ """
+ b = event.current_buffer
+
+ # When already navigating through completions, select the next one.
+ if b.complete_state:
+ b.complete_next()
+ else:
+ b.start_completion(insert_common_part=True)
+
+
+def display_completions_like_readline(event: E) -> None:
+ """
+ Key binding handler for readline-style tab completion.
+ This is meant to be as similar as possible to the way how readline displays
+ completions.
+
+ Generate the completions immediately (blocking) and display them above the
+ prompt in columns.
+
+ Usage::
+
+ # Call this handler when 'Tab' has been pressed.
+ key_bindings.add(Keys.ControlI)(display_completions_like_readline)
+ """
+ # Request completions.
+ b = event.current_buffer
+ if b.completer is None:
+ return
+ complete_event = CompleteEvent(completion_requested=True)
+ completions = list(b.completer.get_completions(b.document, complete_event))
+
+ # Calculate the common suffix.
+ common_suffix = get_common_complete_suffix(b.document, completions)
+
+ # One completion: insert it.
+ if len(completions) == 1:
+ b.delete_before_cursor(-completions[0].start_position)
+ b.insert_text(completions[0].text)
+ # Multiple completions with common part.
+ elif common_suffix:
+ b.insert_text(common_suffix)
+ # Otherwise: display all completions.
+ elif completions:
+ _display_completions_like_readline(event.app, completions)
+
+
+def _display_completions_like_readline(
+ app: Application[object], completions: list[Completion]
+) -> asyncio.Task[None]:
+ """
+ Display the list of completions in columns above the prompt.
+ This will ask for a confirmation if there are too many completions to fit
+ on a single page and provide a paginator to walk through them.
+ """
+ from prompt_toolkit.formatted_text import to_formatted_text
+ from prompt_toolkit.shortcuts.prompt import create_confirm_session
+
+ # Get terminal dimensions.
+ term_size = app.output.get_size()
+ term_width = term_size.columns
+ term_height = term_size.rows
+
+ # Calculate amount of required columns/rows for displaying the
+ # completions. (Keep in mind that completions are displayed
+ # alphabetically column-wise.)
+ max_compl_width = min(
+ term_width, max(get_cwidth(c.display_text) for c in completions) + 1
+ )
+ column_count = max(1, term_width // max_compl_width)
+ completions_per_page = column_count * (term_height - 1)
+ page_count = int(math.ceil(len(completions) / float(completions_per_page)))
+ # Note: math.ceil can return float on Python2.
+
+ def display(page: int) -> None:
+ # Display completions.
+ page_completions = completions[
+ page * completions_per_page : (page + 1) * completions_per_page
+ ]
+
+ page_row_count = int(math.ceil(len(page_completions) / float(column_count)))
+ page_columns = [
+ page_completions[i * page_row_count : (i + 1) * page_row_count]
+ for i in range(column_count)
+ ]
+
+ result: StyleAndTextTuples = []
+
+ for r in range(page_row_count):
+ for c in range(column_count):
+ try:
+ completion = page_columns[c][r]
+ style = "class:readline-like-completions.completion " + (
+ completion.style or ""
+ )
+
+ result.extend(to_formatted_text(completion.display, style=style))
+
+ # Add padding.
+ padding = max_compl_width - get_cwidth(completion.display_text)
+ result.append((completion.style, " " * padding))
+ except IndexError:
+ pass
+ result.append(("", "\n"))
+
+ app.print_text(to_formatted_text(result, "class:readline-like-completions"))
+
+ # User interaction through an application generator function.
+ async def run_compl() -> None:
+ "Coroutine."
+ async with in_terminal(render_cli_done=True):
+ if len(completions) > completions_per_page:
+ # Ask confirmation if it doesn't fit on the screen.
+ confirm = await create_confirm_session(
+ f"Display all {len(completions)} possibilities?",
+ ).prompt_async()
+
+ if confirm:
+ # Display pages.
+ for page in range(page_count):
+ display(page)
+
+ if page != page_count - 1:
+ # Display --MORE-- and go to the next page.
+ show_more = await _create_more_session(
+ "--MORE--"
+ ).prompt_async()
+
+ if not show_more:
+ return
+ else:
+ app.output.flush()
+ else:
+ # Display all completions.
+ display(0)
+
+ return app.create_background_task(run_compl())
+
+
+def _create_more_session(message: str = "--MORE--") -> PromptSession[bool]:
+ """
+ Create a `PromptSession` object for displaying the "--MORE--".
+ """
+ from prompt_toolkit.shortcuts import PromptSession
+
+ bindings = KeyBindings()
+
+ @bindings.add(" ")
+ @bindings.add("y")
+ @bindings.add("Y")
+ @bindings.add(Keys.ControlJ)
+ @bindings.add(Keys.ControlM)
+ @bindings.add(Keys.ControlI) # Tab.
+ def _yes(event: E) -> None:
+ event.app.exit(result=True)
+
+ @bindings.add("n")
+ @bindings.add("N")
+ @bindings.add("q")
+ @bindings.add("Q")
+ @bindings.add(Keys.ControlC)
+ def _no(event: E) -> None:
+ event.app.exit(result=False)
+
+ @bindings.add(Keys.Any)
+ def _ignore(event: E) -> None:
+ "Disable inserting of text."
+
+ return PromptSession(message, key_bindings=bindings, erase_when_done=True)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/cpr.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/cpr.py
new file mode 100644
index 0000000000000000000000000000000000000000..cd9df0a60e725aa46a1906b6614f8f3f48e97fa6
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/cpr.py
@@ -0,0 +1,30 @@
+from __future__ import annotations
+
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+from prompt_toolkit.keys import Keys
+
+from ..key_bindings import KeyBindings
+
+__all__ = [
+ "load_cpr_bindings",
+]
+
+E = KeyPressEvent
+
+
+def load_cpr_bindings() -> KeyBindings:
+ key_bindings = KeyBindings()
+
+ @key_bindings.add(Keys.CPRResponse, save_before=lambda e: False)
+ def _(event: E) -> None:
+ """
+ Handle incoming Cursor-Position-Request response.
+ """
+ # The incoming data looks like u'\x1b[35;1R'
+ # Parse row/col information.
+ row, col = map(int, event.data[2:-1].split(";"))
+
+ # Report absolute cursor position to the renderer.
+ event.app.renderer.report_absolute_cursor_row(row)
+
+ return key_bindings
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/emacs.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/emacs.py
new file mode 100644
index 0000000000000000000000000000000000000000..207afba7ac6d0033e71877794ffcbb276a4e2ba3
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/emacs.py
@@ -0,0 +1,563 @@
+# pylint: disable=function-redefined
+from __future__ import annotations
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.buffer import Buffer, indent, unindent
+from prompt_toolkit.completion import CompleteEvent
+from prompt_toolkit.filters import (
+ Condition,
+ emacs_insert_mode,
+ emacs_mode,
+ has_arg,
+ has_selection,
+ in_paste_mode,
+ is_multiline,
+ is_read_only,
+ shift_selection_mode,
+ vi_search_direction_reversed,
+)
+from prompt_toolkit.key_binding.key_bindings import Binding
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+from prompt_toolkit.keys import Keys
+from prompt_toolkit.selection import SelectionType
+
+from ..key_bindings import ConditionalKeyBindings, KeyBindings, KeyBindingsBase
+from .named_commands import get_by_name
+
+__all__ = [
+ "load_emacs_bindings",
+ "load_emacs_search_bindings",
+ "load_emacs_shift_selection_bindings",
+]
+
+E = KeyPressEvent
+
+
+@Condition
+def is_returnable() -> bool:
+ return get_app().current_buffer.is_returnable
+
+
+@Condition
+def is_arg() -> bool:
+ return get_app().key_processor.arg == "-"
+
+
+def load_emacs_bindings() -> KeyBindingsBase:
+ """
+ Some e-macs extensions.
+ """
+ # Overview of Readline emacs commands:
+ # http://www.catonmat.net/download/readline-emacs-editing-mode-cheat-sheet.pdf
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+
+ insert_mode = emacs_insert_mode
+
+ @handle("escape")
+ def _esc(event: E) -> None:
+ """
+ By default, ignore escape key.
+
+ (If we don't put this here, and Esc is followed by a key which sequence
+ is not handled, we'll insert an Escape character in the input stream.
+ Something we don't want and happens to easily in emacs mode.
+ Further, people can always use ControlQ to do a quoted insert.)
+ """
+ pass
+
+ handle("c-a")(get_by_name("beginning-of-line"))
+ handle("c-b")(get_by_name("backward-char"))
+ handle("c-delete", filter=insert_mode)(get_by_name("kill-word"))
+ handle("c-e")(get_by_name("end-of-line"))
+ handle("c-f")(get_by_name("forward-char"))
+ handle("c-left")(get_by_name("backward-word"))
+ handle("c-right")(get_by_name("forward-word"))
+ handle("c-x", "r", "y", filter=insert_mode)(get_by_name("yank"))
+ handle("c-y", filter=insert_mode)(get_by_name("yank"))
+ handle("escape", "b")(get_by_name("backward-word"))
+ handle("escape", "c", filter=insert_mode)(get_by_name("capitalize-word"))
+ handle("escape", "d", filter=insert_mode)(get_by_name("kill-word"))
+ handle("escape", "f")(get_by_name("forward-word"))
+ handle("escape", "l", filter=insert_mode)(get_by_name("downcase-word"))
+ handle("escape", "u", filter=insert_mode)(get_by_name("uppercase-word"))
+ handle("escape", "y", filter=insert_mode)(get_by_name("yank-pop"))
+ handle("escape", "backspace", filter=insert_mode)(get_by_name("backward-kill-word"))
+ handle("escape", "\\", filter=insert_mode)(get_by_name("delete-horizontal-space"))
+
+ handle("c-home")(get_by_name("beginning-of-buffer"))
+ handle("c-end")(get_by_name("end-of-buffer"))
+
+ handle("c-_", save_before=(lambda e: False), filter=insert_mode)(
+ get_by_name("undo")
+ )
+
+ handle("c-x", "c-u", save_before=(lambda e: False), filter=insert_mode)(
+ get_by_name("undo")
+ )
+
+ handle("escape", "<", filter=~has_selection)(get_by_name("beginning-of-history"))
+ handle("escape", ">", filter=~has_selection)(get_by_name("end-of-history"))
+
+ handle("escape", ".", filter=insert_mode)(get_by_name("yank-last-arg"))
+ handle("escape", "_", filter=insert_mode)(get_by_name("yank-last-arg"))
+ handle("escape", "c-y", filter=insert_mode)(get_by_name("yank-nth-arg"))
+ handle("escape", "#", filter=insert_mode)(get_by_name("insert-comment"))
+ handle("c-o")(get_by_name("operate-and-get-next"))
+
+ # ControlQ does a quoted insert. Not that for vt100 terminals, you have to
+ # disable flow control by running ``stty -ixon``, otherwise Ctrl-Q and
+ # Ctrl-S are captured by the terminal.
+ handle("c-q", filter=~has_selection)(get_by_name("quoted-insert"))
+
+ handle("c-x", "(")(get_by_name("start-kbd-macro"))
+ handle("c-x", ")")(get_by_name("end-kbd-macro"))
+ handle("c-x", "e")(get_by_name("call-last-kbd-macro"))
+
+ @handle("c-n")
+ def _next(event: E) -> None:
+ "Next line."
+ event.current_buffer.auto_down()
+
+ @handle("c-p")
+ def _prev(event: E) -> None:
+ "Previous line."
+ event.current_buffer.auto_up(count=event.arg)
+
+ def handle_digit(c: str) -> None:
+ """
+ Handle input of arguments.
+ The first number needs to be preceded by escape.
+ """
+
+ @handle(c, filter=has_arg)
+ @handle("escape", c)
+ def _(event: E) -> None:
+ event.append_to_arg_count(c)
+
+ for c in "0123456789":
+ handle_digit(c)
+
+ @handle("escape", "-", filter=~has_arg)
+ def _meta_dash(event: E) -> None:
+ """"""
+ if event._arg is None:
+ event.append_to_arg_count("-")
+
+ @handle("-", filter=is_arg)
+ def _dash(event: E) -> None:
+ """
+ When '-' is typed again, after exactly '-' has been given as an
+ argument, ignore this.
+ """
+ event.app.key_processor.arg = "-"
+
+ # Meta + Enter: always accept input.
+ handle("escape", "enter", filter=insert_mode & is_returnable)(
+ get_by_name("accept-line")
+ )
+
+ # Enter: accept input in single line mode.
+ handle("enter", filter=insert_mode & is_returnable & ~is_multiline)(
+ get_by_name("accept-line")
+ )
+
+ def character_search(buff: Buffer, char: str, count: int) -> None:
+ if count < 0:
+ match = buff.document.find_backwards(
+ char, in_current_line=True, count=-count
+ )
+ else:
+ match = buff.document.find(char, in_current_line=True, count=count)
+
+ if match is not None:
+ buff.cursor_position += match
+
+ @handle("c-]", Keys.Any)
+ def _goto_char(event: E) -> None:
+ "When Ctl-] + a character is pressed. go to that character."
+ # Also named 'character-search'
+ character_search(event.current_buffer, event.data, event.arg)
+
+ @handle("escape", "c-]", Keys.Any)
+ def _goto_char_backwards(event: E) -> None:
+ "Like Ctl-], but backwards."
+ # Also named 'character-search-backward'
+ character_search(event.current_buffer, event.data, -event.arg)
+
+ @handle("escape", "a")
+ def _prev_sentence(event: E) -> None:
+ "Previous sentence."
+ # TODO:
+
+ @handle("escape", "e")
+ def _end_of_sentence(event: E) -> None:
+ "Move to end of sentence."
+ # TODO:
+
+ @handle("escape", "t", filter=insert_mode)
+ def _swap_characters(event: E) -> None:
+ """
+ Swap the last two words before the cursor.
+ """
+ # TODO
+
+ @handle("escape", "*", filter=insert_mode)
+ def _insert_all_completions(event: E) -> None:
+ """
+ `meta-*`: Insert all possible completions of the preceding text.
+ """
+ buff = event.current_buffer
+
+ # List all completions.
+ complete_event = CompleteEvent(text_inserted=False, completion_requested=True)
+ completions = list(
+ buff.completer.get_completions(buff.document, complete_event)
+ )
+
+ # Insert them.
+ text_to_insert = " ".join(c.text for c in completions)
+ buff.insert_text(text_to_insert)
+
+ @handle("c-x", "c-x")
+ def _toggle_start_end(event: E) -> None:
+ """
+ Move cursor back and forth between the start and end of the current
+ line.
+ """
+ buffer = event.current_buffer
+
+ if buffer.document.is_cursor_at_the_end_of_line:
+ buffer.cursor_position += buffer.document.get_start_of_line_position(
+ after_whitespace=False
+ )
+ else:
+ buffer.cursor_position += buffer.document.get_end_of_line_position()
+
+ @handle("c-@") # Control-space or Control-@
+ def _start_selection(event: E) -> None:
+ """
+ Start of the selection (if the current buffer is not empty).
+ """
+ # Take the current cursor position as the start of this selection.
+ buff = event.current_buffer
+ if buff.text:
+ buff.start_selection(selection_type=SelectionType.CHARACTERS)
+
+ @handle("c-g", filter=~has_selection)
+ def _cancel(event: E) -> None:
+ """
+ Control + G: Cancel completion menu and validation state.
+ """
+ event.current_buffer.complete_state = None
+ event.current_buffer.validation_error = None
+
+ @handle("c-g", filter=has_selection)
+ def _cancel_selection(event: E) -> None:
+ """
+ Cancel selection.
+ """
+ event.current_buffer.exit_selection()
+
+ @handle("c-w", filter=has_selection)
+ @handle("c-x", "r", "k", filter=has_selection)
+ def _cut(event: E) -> None:
+ """
+ Cut selected text.
+ """
+ data = event.current_buffer.cut_selection()
+ event.app.clipboard.set_data(data)
+
+ @handle("escape", "w", filter=has_selection)
+ def _copy(event: E) -> None:
+ """
+ Copy selected text.
+ """
+ data = event.current_buffer.copy_selection()
+ event.app.clipboard.set_data(data)
+
+ @handle("escape", "left")
+ def _start_of_word(event: E) -> None:
+ """
+ Cursor to start of previous word.
+ """
+ buffer = event.current_buffer
+ buffer.cursor_position += (
+ buffer.document.find_previous_word_beginning(count=event.arg) or 0
+ )
+
+ @handle("escape", "right")
+ def _start_next_word(event: E) -> None:
+ """
+ Cursor to start of next word.
+ """
+ buffer = event.current_buffer
+ buffer.cursor_position += (
+ buffer.document.find_next_word_beginning(count=event.arg)
+ or buffer.document.get_end_of_document_position()
+ )
+
+ @handle("escape", "/", filter=insert_mode)
+ def _complete(event: E) -> None:
+ """
+ M-/: Complete.
+ """
+ b = event.current_buffer
+ if b.complete_state:
+ b.complete_next()
+ else:
+ b.start_completion(select_first=True)
+
+ @handle("c-c", ">", filter=has_selection)
+ def _indent(event: E) -> None:
+ """
+ Indent selected text.
+ """
+ buffer = event.current_buffer
+
+ buffer.cursor_position += buffer.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+
+ from_, to = buffer.document.selection_range()
+ from_, _ = buffer.document.translate_index_to_position(from_)
+ to, _ = buffer.document.translate_index_to_position(to)
+
+ indent(buffer, from_, to + 1, count=event.arg)
+
+ @handle("c-c", "<", filter=has_selection)
+ def _unindent(event: E) -> None:
+ """
+ Unindent selected text.
+ """
+ buffer = event.current_buffer
+
+ from_, to = buffer.document.selection_range()
+ from_, _ = buffer.document.translate_index_to_position(from_)
+ to, _ = buffer.document.translate_index_to_position(to)
+
+ unindent(buffer, from_, to + 1, count=event.arg)
+
+ return ConditionalKeyBindings(key_bindings, emacs_mode)
+
+
+def load_emacs_search_bindings() -> KeyBindingsBase:
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+ from . import search
+
+ # NOTE: We don't bind 'Escape' to 'abort_search'. The reason is that we
+ # want Alt+Enter to accept input directly in incremental search mode.
+ # Instead, we have double escape.
+
+ handle("c-r")(search.start_reverse_incremental_search)
+ handle("c-s")(search.start_forward_incremental_search)
+
+ handle("c-c")(search.abort_search)
+ handle("c-g")(search.abort_search)
+ handle("c-r")(search.reverse_incremental_search)
+ handle("c-s")(search.forward_incremental_search)
+ handle("up")(search.reverse_incremental_search)
+ handle("down")(search.forward_incremental_search)
+ handle("enter")(search.accept_search)
+
+ # Handling of escape.
+ handle("escape", eager=True)(search.accept_search)
+
+ # Like Readline, it's more natural to accept the search when escape has
+ # been pressed, however instead the following two bindings could be used
+ # instead.
+ # #handle('escape', 'escape', eager=True)(search.abort_search)
+ # #handle('escape', 'enter', eager=True)(search.accept_search_and_accept_input)
+
+ # If Read-only: also include the following key bindings:
+
+ # '/' and '?' key bindings for searching, just like Vi mode.
+ handle("?", filter=is_read_only & ~vi_search_direction_reversed)(
+ search.start_reverse_incremental_search
+ )
+ handle("/", filter=is_read_only & ~vi_search_direction_reversed)(
+ search.start_forward_incremental_search
+ )
+ handle("?", filter=is_read_only & vi_search_direction_reversed)(
+ search.start_forward_incremental_search
+ )
+ handle("/", filter=is_read_only & vi_search_direction_reversed)(
+ search.start_reverse_incremental_search
+ )
+
+ @handle("n", filter=is_read_only)
+ def _jump_next(event: E) -> None:
+ "Jump to next match."
+ event.current_buffer.apply_search(
+ event.app.current_search_state,
+ include_current_position=False,
+ count=event.arg,
+ )
+
+ @handle("N", filter=is_read_only)
+ def _jump_prev(event: E) -> None:
+ "Jump to previous match."
+ event.current_buffer.apply_search(
+ ~event.app.current_search_state,
+ include_current_position=False,
+ count=event.arg,
+ )
+
+ return ConditionalKeyBindings(key_bindings, emacs_mode)
+
+
+def load_emacs_shift_selection_bindings() -> KeyBindingsBase:
+ """
+ Bindings to select text with shift + cursor movements
+ """
+
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+
+ def unshift_move(event: E) -> None:
+ """
+ Used for the shift selection mode. When called with
+ a shift + movement key press event, moves the cursor
+ as if shift is not pressed.
+ """
+ key = event.key_sequence[0].key
+
+ if key == Keys.ShiftUp:
+ event.current_buffer.auto_up(count=event.arg)
+ return
+ if key == Keys.ShiftDown:
+ event.current_buffer.auto_down(count=event.arg)
+ return
+
+ # the other keys are handled through their readline command
+ key_to_command: dict[Keys | str, str] = {
+ Keys.ShiftLeft: "backward-char",
+ Keys.ShiftRight: "forward-char",
+ Keys.ShiftHome: "beginning-of-line",
+ Keys.ShiftEnd: "end-of-line",
+ Keys.ControlShiftLeft: "backward-word",
+ Keys.ControlShiftRight: "forward-word",
+ Keys.ControlShiftHome: "beginning-of-buffer",
+ Keys.ControlShiftEnd: "end-of-buffer",
+ }
+
+ try:
+ # Both the dict lookup and `get_by_name` can raise KeyError.
+ binding = get_by_name(key_to_command[key])
+ except KeyError:
+ pass
+ else: # (`else` is not really needed here.)
+ if isinstance(binding, Binding):
+ # (It should always be a binding here)
+ binding.call(event)
+
+ @handle("s-left", filter=~has_selection)
+ @handle("s-right", filter=~has_selection)
+ @handle("s-up", filter=~has_selection)
+ @handle("s-down", filter=~has_selection)
+ @handle("s-home", filter=~has_selection)
+ @handle("s-end", filter=~has_selection)
+ @handle("c-s-left", filter=~has_selection)
+ @handle("c-s-right", filter=~has_selection)
+ @handle("c-s-home", filter=~has_selection)
+ @handle("c-s-end", filter=~has_selection)
+ def _start_selection(event: E) -> None:
+ """
+ Start selection with shift + movement.
+ """
+ # Take the current cursor position as the start of this selection.
+ buff = event.current_buffer
+ if buff.text:
+ buff.start_selection(selection_type=SelectionType.CHARACTERS)
+
+ if buff.selection_state is not None:
+ # (`selection_state` should never be `None`, it is created by
+ # `start_selection`.)
+ buff.selection_state.enter_shift_mode()
+
+ # Then move the cursor
+ original_position = buff.cursor_position
+ unshift_move(event)
+ if buff.cursor_position == original_position:
+ # Cursor didn't actually move - so cancel selection
+ # to avoid having an empty selection
+ buff.exit_selection()
+
+ @handle("s-left", filter=shift_selection_mode)
+ @handle("s-right", filter=shift_selection_mode)
+ @handle("s-up", filter=shift_selection_mode)
+ @handle("s-down", filter=shift_selection_mode)
+ @handle("s-home", filter=shift_selection_mode)
+ @handle("s-end", filter=shift_selection_mode)
+ @handle("c-s-left", filter=shift_selection_mode)
+ @handle("c-s-right", filter=shift_selection_mode)
+ @handle("c-s-home", filter=shift_selection_mode)
+ @handle("c-s-end", filter=shift_selection_mode)
+ def _extend_selection(event: E) -> None:
+ """
+ Extend the selection
+ """
+ # Just move the cursor, like shift was not pressed
+ unshift_move(event)
+ buff = event.current_buffer
+
+ if buff.selection_state is not None:
+ if buff.cursor_position == buff.selection_state.original_cursor_position:
+ # selection is now empty, so cancel selection
+ buff.exit_selection()
+
+ @handle(Keys.Any, filter=shift_selection_mode)
+ def _replace_selection(event: E) -> None:
+ """
+ Replace selection by what is typed
+ """
+ event.current_buffer.cut_selection()
+ get_by_name("self-insert").call(event)
+
+ @handle("enter", filter=shift_selection_mode & is_multiline)
+ def _newline(event: E) -> None:
+ """
+ A newline replaces the selection
+ """
+ event.current_buffer.cut_selection()
+ event.current_buffer.newline(copy_margin=not in_paste_mode())
+
+ @handle("backspace", filter=shift_selection_mode)
+ def _delete(event: E) -> None:
+ """
+ Delete selection.
+ """
+ event.current_buffer.cut_selection()
+
+ @handle("c-y", filter=shift_selection_mode)
+ def _yank(event: E) -> None:
+ """
+ In shift selection mode, yanking (pasting) replace the selection.
+ """
+ buff = event.current_buffer
+ if buff.selection_state:
+ buff.cut_selection()
+ get_by_name("yank").call(event)
+
+ # moving the cursor in shift selection mode cancels the selection
+ @handle("left", filter=shift_selection_mode)
+ @handle("right", filter=shift_selection_mode)
+ @handle("up", filter=shift_selection_mode)
+ @handle("down", filter=shift_selection_mode)
+ @handle("home", filter=shift_selection_mode)
+ @handle("end", filter=shift_selection_mode)
+ @handle("c-left", filter=shift_selection_mode)
+ @handle("c-right", filter=shift_selection_mode)
+ @handle("c-home", filter=shift_selection_mode)
+ @handle("c-end", filter=shift_selection_mode)
+ def _cancel(event: E) -> None:
+ """
+ Cancel selection.
+ """
+ event.current_buffer.exit_selection()
+ # we then process the cursor movement
+ key_press = event.key_sequence[0]
+ event.key_processor.feed(key_press, first=True)
+
+ return ConditionalKeyBindings(key_bindings, emacs_mode)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/focus.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/focus.py
new file mode 100644
index 0000000000000000000000000000000000000000..24aa3ce39ba798d47668212fa1c6a7afed069032
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/focus.py
@@ -0,0 +1,26 @@
+from __future__ import annotations
+
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+
+__all__ = [
+ "focus_next",
+ "focus_previous",
+]
+
+E = KeyPressEvent
+
+
+def focus_next(event: E) -> None:
+ """
+ Focus the next visible Window.
+ (Often bound to the `Tab` key.)
+ """
+ event.app.layout.focus_next()
+
+
+def focus_previous(event: E) -> None:
+ """
+ Focus the previous visible Window.
+ (Often bound to the `BackTab` key.)
+ """
+ event.app.layout.focus_previous()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/mouse.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/mouse.py
new file mode 100644
index 0000000000000000000000000000000000000000..cb426ce7bc8dee34ae7db425c0de526795599ac9
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/mouse.py
@@ -0,0 +1,348 @@
+from __future__ import annotations
+
+import sys
+from typing import TYPE_CHECKING
+
+from prompt_toolkit.data_structures import Point
+from prompt_toolkit.key_binding.key_processor import KeyPress, KeyPressEvent
+from prompt_toolkit.keys import Keys
+from prompt_toolkit.mouse_events import (
+ MouseButton,
+ MouseEvent,
+ MouseEventType,
+ MouseModifier,
+)
+
+from ..key_bindings import KeyBindings
+
+if TYPE_CHECKING:
+ from prompt_toolkit.key_binding.key_bindings import NotImplementedOrNone
+
+__all__ = [
+ "load_mouse_bindings",
+]
+
+E = KeyPressEvent
+
+# fmt: off
+SCROLL_UP = MouseEventType.SCROLL_UP
+SCROLL_DOWN = MouseEventType.SCROLL_DOWN
+MOUSE_DOWN = MouseEventType.MOUSE_DOWN
+MOUSE_MOVE = MouseEventType.MOUSE_MOVE
+MOUSE_UP = MouseEventType.MOUSE_UP
+
+NO_MODIFIER : frozenset[MouseModifier] = frozenset()
+SHIFT : frozenset[MouseModifier] = frozenset({MouseModifier.SHIFT})
+ALT : frozenset[MouseModifier] = frozenset({MouseModifier.ALT})
+SHIFT_ALT : frozenset[MouseModifier] = frozenset({MouseModifier.SHIFT, MouseModifier.ALT})
+CONTROL : frozenset[MouseModifier] = frozenset({MouseModifier.CONTROL})
+SHIFT_CONTROL : frozenset[MouseModifier] = frozenset({MouseModifier.SHIFT, MouseModifier.CONTROL})
+ALT_CONTROL : frozenset[MouseModifier] = frozenset({MouseModifier.ALT, MouseModifier.CONTROL})
+SHIFT_ALT_CONTROL: frozenset[MouseModifier] = frozenset({MouseModifier.SHIFT, MouseModifier.ALT, MouseModifier.CONTROL})
+UNKNOWN_MODIFIER : frozenset[MouseModifier] = frozenset()
+
+LEFT = MouseButton.LEFT
+MIDDLE = MouseButton.MIDDLE
+RIGHT = MouseButton.RIGHT
+NO_BUTTON = MouseButton.NONE
+UNKNOWN_BUTTON = MouseButton.UNKNOWN
+
+xterm_sgr_mouse_events = {
+ ( 0, "m") : (LEFT, MOUSE_UP, NO_MODIFIER), # left_up 0+ + + =0
+ ( 4, "m") : (LEFT, MOUSE_UP, SHIFT), # left_up Shift 0+4+ + =4
+ ( 8, "m") : (LEFT, MOUSE_UP, ALT), # left_up Alt 0+ +8+ =8
+ (12, "m") : (LEFT, MOUSE_UP, SHIFT_ALT), # left_up Shift Alt 0+4+8+ =12
+ (16, "m") : (LEFT, MOUSE_UP, CONTROL), # left_up Control 0+ + +16=16
+ (20, "m") : (LEFT, MOUSE_UP, SHIFT_CONTROL), # left_up Shift Control 0+4+ +16=20
+ (24, "m") : (LEFT, MOUSE_UP, ALT_CONTROL), # left_up Alt Control 0+ +8+16=24
+ (28, "m") : (LEFT, MOUSE_UP, SHIFT_ALT_CONTROL), # left_up Shift Alt Control 0+4+8+16=28
+
+ ( 1, "m") : (MIDDLE, MOUSE_UP, NO_MODIFIER), # middle_up 1+ + + =1
+ ( 5, "m") : (MIDDLE, MOUSE_UP, SHIFT), # middle_up Shift 1+4+ + =5
+ ( 9, "m") : (MIDDLE, MOUSE_UP, ALT), # middle_up Alt 1+ +8+ =9
+ (13, "m") : (MIDDLE, MOUSE_UP, SHIFT_ALT), # middle_up Shift Alt 1+4+8+ =13
+ (17, "m") : (MIDDLE, MOUSE_UP, CONTROL), # middle_up Control 1+ + +16=17
+ (21, "m") : (MIDDLE, MOUSE_UP, SHIFT_CONTROL), # middle_up Shift Control 1+4+ +16=21
+ (25, "m") : (MIDDLE, MOUSE_UP, ALT_CONTROL), # middle_up Alt Control 1+ +8+16=25
+ (29, "m") : (MIDDLE, MOUSE_UP, SHIFT_ALT_CONTROL), # middle_up Shift Alt Control 1+4+8+16=29
+
+ ( 2, "m") : (RIGHT, MOUSE_UP, NO_MODIFIER), # right_up 2+ + + =2
+ ( 6, "m") : (RIGHT, MOUSE_UP, SHIFT), # right_up Shift 2+4+ + =6
+ (10, "m") : (RIGHT, MOUSE_UP, ALT), # right_up Alt 2+ +8+ =10
+ (14, "m") : (RIGHT, MOUSE_UP, SHIFT_ALT), # right_up Shift Alt 2+4+8+ =14
+ (18, "m") : (RIGHT, MOUSE_UP, CONTROL), # right_up Control 2+ + +16=18
+ (22, "m") : (RIGHT, MOUSE_UP, SHIFT_CONTROL), # right_up Shift Control 2+4+ +16=22
+ (26, "m") : (RIGHT, MOUSE_UP, ALT_CONTROL), # right_up Alt Control 2+ +8+16=26
+ (30, "m") : (RIGHT, MOUSE_UP, SHIFT_ALT_CONTROL), # right_up Shift Alt Control 2+4+8+16=30
+
+ ( 0, "M") : (LEFT, MOUSE_DOWN, NO_MODIFIER), # left_down 0+ + + =0
+ ( 4, "M") : (LEFT, MOUSE_DOWN, SHIFT), # left_down Shift 0+4+ + =4
+ ( 8, "M") : (LEFT, MOUSE_DOWN, ALT), # left_down Alt 0+ +8+ =8
+ (12, "M") : (LEFT, MOUSE_DOWN, SHIFT_ALT), # left_down Shift Alt 0+4+8+ =12
+ (16, "M") : (LEFT, MOUSE_DOWN, CONTROL), # left_down Control 0+ + +16=16
+ (20, "M") : (LEFT, MOUSE_DOWN, SHIFT_CONTROL), # left_down Shift Control 0+4+ +16=20
+ (24, "M") : (LEFT, MOUSE_DOWN, ALT_CONTROL), # left_down Alt Control 0+ +8+16=24
+ (28, "M") : (LEFT, MOUSE_DOWN, SHIFT_ALT_CONTROL), # left_down Shift Alt Control 0+4+8+16=28
+
+ ( 1, "M") : (MIDDLE, MOUSE_DOWN, NO_MODIFIER), # middle_down 1+ + + =1
+ ( 5, "M") : (MIDDLE, MOUSE_DOWN, SHIFT), # middle_down Shift 1+4+ + =5
+ ( 9, "M") : (MIDDLE, MOUSE_DOWN, ALT), # middle_down Alt 1+ +8+ =9
+ (13, "M") : (MIDDLE, MOUSE_DOWN, SHIFT_ALT), # middle_down Shift Alt 1+4+8+ =13
+ (17, "M") : (MIDDLE, MOUSE_DOWN, CONTROL), # middle_down Control 1+ + +16=17
+ (21, "M") : (MIDDLE, MOUSE_DOWN, SHIFT_CONTROL), # middle_down Shift Control 1+4+ +16=21
+ (25, "M") : (MIDDLE, MOUSE_DOWN, ALT_CONTROL), # middle_down Alt Control 1+ +8+16=25
+ (29, "M") : (MIDDLE, MOUSE_DOWN, SHIFT_ALT_CONTROL), # middle_down Shift Alt Control 1+4+8+16=29
+
+ ( 2, "M") : (RIGHT, MOUSE_DOWN, NO_MODIFIER), # right_down 2+ + + =2
+ ( 6, "M") : (RIGHT, MOUSE_DOWN, SHIFT), # right_down Shift 2+4+ + =6
+ (10, "M") : (RIGHT, MOUSE_DOWN, ALT), # right_down Alt 2+ +8+ =10
+ (14, "M") : (RIGHT, MOUSE_DOWN, SHIFT_ALT), # right_down Shift Alt 2+4+8+ =14
+ (18, "M") : (RIGHT, MOUSE_DOWN, CONTROL), # right_down Control 2+ + +16=18
+ (22, "M") : (RIGHT, MOUSE_DOWN, SHIFT_CONTROL), # right_down Shift Control 2+4+ +16=22
+ (26, "M") : (RIGHT, MOUSE_DOWN, ALT_CONTROL), # right_down Alt Control 2+ +8+16=26
+ (30, "M") : (RIGHT, MOUSE_DOWN, SHIFT_ALT_CONTROL), # right_down Shift Alt Control 2+4+8+16=30
+
+ (32, "M") : (LEFT, MOUSE_MOVE, NO_MODIFIER), # left_drag 32+ + + =32
+ (36, "M") : (LEFT, MOUSE_MOVE, SHIFT), # left_drag Shift 32+4+ + =36
+ (40, "M") : (LEFT, MOUSE_MOVE, ALT), # left_drag Alt 32+ +8+ =40
+ (44, "M") : (LEFT, MOUSE_MOVE, SHIFT_ALT), # left_drag Shift Alt 32+4+8+ =44
+ (48, "M") : (LEFT, MOUSE_MOVE, CONTROL), # left_drag Control 32+ + +16=48
+ (52, "M") : (LEFT, MOUSE_MOVE, SHIFT_CONTROL), # left_drag Shift Control 32+4+ +16=52
+ (56, "M") : (LEFT, MOUSE_MOVE, ALT_CONTROL), # left_drag Alt Control 32+ +8+16=56
+ (60, "M") : (LEFT, MOUSE_MOVE, SHIFT_ALT_CONTROL), # left_drag Shift Alt Control 32+4+8+16=60
+
+ (33, "M") : (MIDDLE, MOUSE_MOVE, NO_MODIFIER), # middle_drag 33+ + + =33
+ (37, "M") : (MIDDLE, MOUSE_MOVE, SHIFT), # middle_drag Shift 33+4+ + =37
+ (41, "M") : (MIDDLE, MOUSE_MOVE, ALT), # middle_drag Alt 33+ +8+ =41
+ (45, "M") : (MIDDLE, MOUSE_MOVE, SHIFT_ALT), # middle_drag Shift Alt 33+4+8+ =45
+ (49, "M") : (MIDDLE, MOUSE_MOVE, CONTROL), # middle_drag Control 33+ + +16=49
+ (53, "M") : (MIDDLE, MOUSE_MOVE, SHIFT_CONTROL), # middle_drag Shift Control 33+4+ +16=53
+ (57, "M") : (MIDDLE, MOUSE_MOVE, ALT_CONTROL), # middle_drag Alt Control 33+ +8+16=57
+ (61, "M") : (MIDDLE, MOUSE_MOVE, SHIFT_ALT_CONTROL), # middle_drag Shift Alt Control 33+4+8+16=61
+
+ (34, "M") : (RIGHT, MOUSE_MOVE, NO_MODIFIER), # right_drag 34+ + + =34
+ (38, "M") : (RIGHT, MOUSE_MOVE, SHIFT), # right_drag Shift 34+4+ + =38
+ (42, "M") : (RIGHT, MOUSE_MOVE, ALT), # right_drag Alt 34+ +8+ =42
+ (46, "M") : (RIGHT, MOUSE_MOVE, SHIFT_ALT), # right_drag Shift Alt 34+4+8+ =46
+ (50, "M") : (RIGHT, MOUSE_MOVE, CONTROL), # right_drag Control 34+ + +16=50
+ (54, "M") : (RIGHT, MOUSE_MOVE, SHIFT_CONTROL), # right_drag Shift Control 34+4+ +16=54
+ (58, "M") : (RIGHT, MOUSE_MOVE, ALT_CONTROL), # right_drag Alt Control 34+ +8+16=58
+ (62, "M") : (RIGHT, MOUSE_MOVE, SHIFT_ALT_CONTROL), # right_drag Shift Alt Control 34+4+8+16=62
+
+ (35, "M") : (NO_BUTTON, MOUSE_MOVE, NO_MODIFIER), # none_drag 35+ + + =35
+ (39, "M") : (NO_BUTTON, MOUSE_MOVE, SHIFT), # none_drag Shift 35+4+ + =39
+ (43, "M") : (NO_BUTTON, MOUSE_MOVE, ALT), # none_drag Alt 35+ +8+ =43
+ (47, "M") : (NO_BUTTON, MOUSE_MOVE, SHIFT_ALT), # none_drag Shift Alt 35+4+8+ =47
+ (51, "M") : (NO_BUTTON, MOUSE_MOVE, CONTROL), # none_drag Control 35+ + +16=51
+ (55, "M") : (NO_BUTTON, MOUSE_MOVE, SHIFT_CONTROL), # none_drag Shift Control 35+4+ +16=55
+ (59, "M") : (NO_BUTTON, MOUSE_MOVE, ALT_CONTROL), # none_drag Alt Control 35+ +8+16=59
+ (63, "M") : (NO_BUTTON, MOUSE_MOVE, SHIFT_ALT_CONTROL), # none_drag Shift Alt Control 35+4+8+16=63
+
+ (64, "M") : (NO_BUTTON, SCROLL_UP, NO_MODIFIER), # scroll_up 64+ + + =64
+ (68, "M") : (NO_BUTTON, SCROLL_UP, SHIFT), # scroll_up Shift 64+4+ + =68
+ (72, "M") : (NO_BUTTON, SCROLL_UP, ALT), # scroll_up Alt 64+ +8+ =72
+ (76, "M") : (NO_BUTTON, SCROLL_UP, SHIFT_ALT), # scroll_up Shift Alt 64+4+8+ =76
+ (80, "M") : (NO_BUTTON, SCROLL_UP, CONTROL), # scroll_up Control 64+ + +16=80
+ (84, "M") : (NO_BUTTON, SCROLL_UP, SHIFT_CONTROL), # scroll_up Shift Control 64+4+ +16=84
+ (88, "M") : (NO_BUTTON, SCROLL_UP, ALT_CONTROL), # scroll_up Alt Control 64+ +8+16=88
+ (92, "M") : (NO_BUTTON, SCROLL_UP, SHIFT_ALT_CONTROL), # scroll_up Shift Alt Control 64+4+8+16=92
+
+ (65, "M") : (NO_BUTTON, SCROLL_DOWN, NO_MODIFIER), # scroll_down 64+ + + =65
+ (69, "M") : (NO_BUTTON, SCROLL_DOWN, SHIFT), # scroll_down Shift 64+4+ + =69
+ (73, "M") : (NO_BUTTON, SCROLL_DOWN, ALT), # scroll_down Alt 64+ +8+ =73
+ (77, "M") : (NO_BUTTON, SCROLL_DOWN, SHIFT_ALT), # scroll_down Shift Alt 64+4+8+ =77
+ (81, "M") : (NO_BUTTON, SCROLL_DOWN, CONTROL), # scroll_down Control 64+ + +16=81
+ (85, "M") : (NO_BUTTON, SCROLL_DOWN, SHIFT_CONTROL), # scroll_down Shift Control 64+4+ +16=85
+ (89, "M") : (NO_BUTTON, SCROLL_DOWN, ALT_CONTROL), # scroll_down Alt Control 64+ +8+16=89
+ (93, "M") : (NO_BUTTON, SCROLL_DOWN, SHIFT_ALT_CONTROL), # scroll_down Shift Alt Control 64+4+8+16=93
+}
+
+typical_mouse_events = {
+ 32: (LEFT , MOUSE_DOWN , UNKNOWN_MODIFIER),
+ 33: (MIDDLE , MOUSE_DOWN , UNKNOWN_MODIFIER),
+ 34: (RIGHT , MOUSE_DOWN , UNKNOWN_MODIFIER),
+ 35: (UNKNOWN_BUTTON , MOUSE_UP , UNKNOWN_MODIFIER),
+
+ 64: (LEFT , MOUSE_MOVE , UNKNOWN_MODIFIER),
+ 65: (MIDDLE , MOUSE_MOVE , UNKNOWN_MODIFIER),
+ 66: (RIGHT , MOUSE_MOVE , UNKNOWN_MODIFIER),
+ 67: (NO_BUTTON , MOUSE_MOVE , UNKNOWN_MODIFIER),
+
+ 96: (NO_BUTTON , SCROLL_UP , UNKNOWN_MODIFIER),
+ 97: (NO_BUTTON , SCROLL_DOWN, UNKNOWN_MODIFIER),
+}
+
+urxvt_mouse_events={
+ 32: (UNKNOWN_BUTTON, MOUSE_DOWN , UNKNOWN_MODIFIER),
+ 35: (UNKNOWN_BUTTON, MOUSE_UP , UNKNOWN_MODIFIER),
+ 96: (NO_BUTTON , SCROLL_UP , UNKNOWN_MODIFIER),
+ 97: (NO_BUTTON , SCROLL_DOWN, UNKNOWN_MODIFIER),
+}
+# fmt:on
+
+
+def load_mouse_bindings() -> KeyBindings:
+ """
+ Key bindings, required for mouse support.
+ (Mouse events enter through the key binding system.)
+ """
+ key_bindings = KeyBindings()
+
+ @key_bindings.add(Keys.Vt100MouseEvent)
+ def _(event: E) -> NotImplementedOrNone:
+ """
+ Handling of incoming mouse event.
+ """
+ # TypicaL: "eSC[MaB*"
+ # Urxvt: "Esc[96;14;13M"
+ # Xterm SGR: "Esc[<64;85;12M"
+
+ # Parse incoming packet.
+ if event.data[2] == "M":
+ # Typical.
+ mouse_event, x, y = map(ord, event.data[3:])
+
+ # TODO: Is it possible to add modifiers here?
+ mouse_button, mouse_event_type, mouse_modifiers = typical_mouse_events[
+ mouse_event
+ ]
+
+ # Handle situations where `PosixStdinReader` used surrogateescapes.
+ if x >= 0xDC00:
+ x -= 0xDC00
+ if y >= 0xDC00:
+ y -= 0xDC00
+
+ x -= 32
+ y -= 32
+ else:
+ # Urxvt and Xterm SGR.
+ # When the '<' is not present, we are not using the Xterm SGR mode,
+ # but Urxvt instead.
+ data = event.data[2:]
+ if data[:1] == "<":
+ sgr = True
+ data = data[1:]
+ else:
+ sgr = False
+
+ # Extract coordinates.
+ mouse_event, x, y = map(int, data[:-1].split(";"))
+ m = data[-1]
+
+ # Parse event type.
+ if sgr:
+ try:
+ (
+ mouse_button,
+ mouse_event_type,
+ mouse_modifiers,
+ ) = xterm_sgr_mouse_events[mouse_event, m]
+ except KeyError:
+ return NotImplemented
+
+ else:
+ # Some other terminals, like urxvt, Hyper terminal, ...
+ (
+ mouse_button,
+ mouse_event_type,
+ mouse_modifiers,
+ ) = urxvt_mouse_events.get(
+ mouse_event, (UNKNOWN_BUTTON, MOUSE_MOVE, UNKNOWN_MODIFIER)
+ )
+
+ x -= 1
+ y -= 1
+
+ # Only handle mouse events when we know the window height.
+ if event.app.renderer.height_is_known and mouse_event_type is not None:
+ # Take region above the layout into account. The reported
+ # coordinates are absolute to the visible part of the terminal.
+ from prompt_toolkit.renderer import HeightIsUnknownError
+
+ try:
+ y -= event.app.renderer.rows_above_layout
+ except HeightIsUnknownError:
+ return NotImplemented
+
+ # Call the mouse handler from the renderer.
+
+ # Note: This can return `NotImplemented` if no mouse handler was
+ # found for this position, or if no repainting needs to
+ # happen. this way, we avoid excessive repaints during mouse
+ # movements.
+ handler = event.app.renderer.mouse_handlers.mouse_handlers[y][x]
+ return handler(
+ MouseEvent(
+ position=Point(x=x, y=y),
+ event_type=mouse_event_type,
+ button=mouse_button,
+ modifiers=mouse_modifiers,
+ )
+ )
+
+ return NotImplemented
+
+ @key_bindings.add(Keys.ScrollUp)
+ def _scroll_up(event: E) -> None:
+ """
+ Scroll up event without cursor position.
+ """
+ # We don't receive a cursor position, so we don't know which window to
+ # scroll. Just send an 'up' key press instead.
+ event.key_processor.feed(KeyPress(Keys.Up), first=True)
+
+ @key_bindings.add(Keys.ScrollDown)
+ def _scroll_down(event: E) -> None:
+ """
+ Scroll down event without cursor position.
+ """
+ event.key_processor.feed(KeyPress(Keys.Down), first=True)
+
+ @key_bindings.add(Keys.WindowsMouseEvent)
+ def _mouse(event: E) -> NotImplementedOrNone:
+ """
+ Handling of mouse events for Windows.
+ """
+ # This key binding should only exist for Windows.
+ if sys.platform == "win32":
+ # Parse data.
+ pieces = event.data.split(";")
+
+ button = MouseButton(pieces[0])
+ event_type = MouseEventType(pieces[1])
+ x = int(pieces[2])
+ y = int(pieces[3])
+
+ # Make coordinates absolute to the visible part of the terminal.
+ output = event.app.renderer.output
+
+ from prompt_toolkit.output.win32 import Win32Output
+ from prompt_toolkit.output.windows10 import Windows10_Output
+
+ if isinstance(output, (Win32Output, Windows10_Output)):
+ screen_buffer_info = output.get_win32_screen_buffer_info()
+ rows_above_cursor = (
+ screen_buffer_info.dwCursorPosition.Y
+ - event.app.renderer._cursor_pos.y
+ )
+ y -= rows_above_cursor
+
+ # Call the mouse event handler.
+ # (Can return `NotImplemented`.)
+ handler = event.app.renderer.mouse_handlers.mouse_handlers[y][x]
+
+ return handler(
+ MouseEvent(
+ position=Point(x=x, y=y),
+ event_type=event_type,
+ button=button,
+ modifiers=UNKNOWN_MODIFIER,
+ )
+ )
+
+ # No mouse handler found. Return `NotImplemented` so that we don't
+ # invalidate the UI.
+ return NotImplemented
+
+ return key_bindings
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/named_commands.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/named_commands.py
new file mode 100644
index 0000000000000000000000000000000000000000..8ea8dd03de4b5ec3869982dbd9b0a167eb0d738e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/named_commands.py
@@ -0,0 +1,691 @@
+"""
+Key bindings which are also known by GNU Readline by the given names.
+
+See: http://www.delorie.com/gnu/docs/readline/rlman_13.html
+"""
+
+from __future__ import annotations
+
+from typing import Callable, TypeVar, Union, cast
+
+from prompt_toolkit.document import Document
+from prompt_toolkit.enums import EditingMode
+from prompt_toolkit.key_binding.key_bindings import Binding, key_binding
+from prompt_toolkit.key_binding.key_processor import KeyPress, KeyPressEvent
+from prompt_toolkit.keys import Keys
+from prompt_toolkit.layout.controls import BufferControl
+from prompt_toolkit.search import SearchDirection
+from prompt_toolkit.selection import PasteMode
+
+from .completion import display_completions_like_readline, generate_completions
+
+__all__ = [
+ "get_by_name",
+]
+
+
+# Typing.
+_Handler = Callable[[KeyPressEvent], None]
+_HandlerOrBinding = Union[_Handler, Binding]
+_T = TypeVar("_T", bound=_HandlerOrBinding)
+E = KeyPressEvent
+
+
+# Registry that maps the Readline command names to their handlers.
+_readline_commands: dict[str, Binding] = {}
+
+
+def register(name: str) -> Callable[[_T], _T]:
+ """
+ Store handler in the `_readline_commands` dictionary.
+ """
+
+ def decorator(handler: _T) -> _T:
+ "`handler` is a callable or Binding."
+ if isinstance(handler, Binding):
+ _readline_commands[name] = handler
+ else:
+ _readline_commands[name] = key_binding()(cast(_Handler, handler))
+
+ return handler
+
+ return decorator
+
+
+def get_by_name(name: str) -> Binding:
+ """
+ Return the handler for the (Readline) command with the given name.
+ """
+ try:
+ return _readline_commands[name]
+ except KeyError as e:
+ raise KeyError(f"Unknown Readline command: {name!r}") from e
+
+
+#
+# Commands for moving
+# See: http://www.delorie.com/gnu/docs/readline/rlman_14.html
+#
+
+
+@register("beginning-of-buffer")
+def beginning_of_buffer(event: E) -> None:
+ """
+ Move to the start of the buffer.
+ """
+ buff = event.current_buffer
+ buff.cursor_position = 0
+
+
+@register("end-of-buffer")
+def end_of_buffer(event: E) -> None:
+ """
+ Move to the end of the buffer.
+ """
+ buff = event.current_buffer
+ buff.cursor_position = len(buff.text)
+
+
+@register("beginning-of-line")
+def beginning_of_line(event: E) -> None:
+ """
+ Move to the start of the current line.
+ """
+ buff = event.current_buffer
+ buff.cursor_position += buff.document.get_start_of_line_position(
+ after_whitespace=False
+ )
+
+
+@register("end-of-line")
+def end_of_line(event: E) -> None:
+ """
+ Move to the end of the line.
+ """
+ buff = event.current_buffer
+ buff.cursor_position += buff.document.get_end_of_line_position()
+
+
+@register("forward-char")
+def forward_char(event: E) -> None:
+ """
+ Move forward a character.
+ """
+ buff = event.current_buffer
+ buff.cursor_position += buff.document.get_cursor_right_position(count=event.arg)
+
+
+@register("backward-char")
+def backward_char(event: E) -> None:
+ "Move back a character."
+ buff = event.current_buffer
+ buff.cursor_position += buff.document.get_cursor_left_position(count=event.arg)
+
+
+@register("forward-word")
+def forward_word(event: E) -> None:
+ """
+ Move forward to the end of the next word. Words are composed of letters and
+ digits.
+ """
+ buff = event.current_buffer
+ pos = buff.document.find_next_word_ending(count=event.arg)
+
+ if pos:
+ buff.cursor_position += pos
+
+
+@register("backward-word")
+def backward_word(event: E) -> None:
+ """
+ Move back to the start of the current or previous word. Words are composed
+ of letters and digits.
+ """
+ buff = event.current_buffer
+ pos = buff.document.find_previous_word_beginning(count=event.arg)
+
+ if pos:
+ buff.cursor_position += pos
+
+
+@register("clear-screen")
+def clear_screen(event: E) -> None:
+ """
+ Clear the screen and redraw everything at the top of the screen.
+ """
+ event.app.renderer.clear()
+
+
+@register("redraw-current-line")
+def redraw_current_line(event: E) -> None:
+ """
+ Refresh the current line.
+ (Readline defines this command, but prompt-toolkit doesn't have it.)
+ """
+ pass
+
+
+#
+# Commands for manipulating the history.
+# See: http://www.delorie.com/gnu/docs/readline/rlman_15.html
+#
+
+
+@register("accept-line")
+def accept_line(event: E) -> None:
+ """
+ Accept the line regardless of where the cursor is.
+ """
+ event.current_buffer.validate_and_handle()
+
+
+@register("previous-history")
+def previous_history(event: E) -> None:
+ """
+ Move `back` through the history list, fetching the previous command.
+ """
+ event.current_buffer.history_backward(count=event.arg)
+
+
+@register("next-history")
+def next_history(event: E) -> None:
+ """
+ Move `forward` through the history list, fetching the next command.
+ """
+ event.current_buffer.history_forward(count=event.arg)
+
+
+@register("beginning-of-history")
+def beginning_of_history(event: E) -> None:
+ """
+ Move to the first line in the history.
+ """
+ event.current_buffer.go_to_history(0)
+
+
+@register("end-of-history")
+def end_of_history(event: E) -> None:
+ """
+ Move to the end of the input history, i.e., the line currently being entered.
+ """
+ event.current_buffer.history_forward(count=10**100)
+ buff = event.current_buffer
+ buff.go_to_history(len(buff._working_lines) - 1)
+
+
+@register("reverse-search-history")
+def reverse_search_history(event: E) -> None:
+ """
+ Search backward starting at the current line and moving `up` through
+ the history as necessary. This is an incremental search.
+ """
+ control = event.app.layout.current_control
+
+ if isinstance(control, BufferControl) and control.search_buffer_control:
+ event.app.current_search_state.direction = SearchDirection.BACKWARD
+ event.app.layout.current_control = control.search_buffer_control
+
+
+#
+# Commands for changing text
+#
+
+
+@register("end-of-file")
+def end_of_file(event: E) -> None:
+ """
+ Exit.
+ """
+ event.app.exit()
+
+
+@register("delete-char")
+def delete_char(event: E) -> None:
+ """
+ Delete character before the cursor.
+ """
+ deleted = event.current_buffer.delete(count=event.arg)
+ if not deleted:
+ event.app.output.bell()
+
+
+@register("backward-delete-char")
+def backward_delete_char(event: E) -> None:
+ """
+ Delete the character behind the cursor.
+ """
+ if event.arg < 0:
+ # When a negative argument has been given, this should delete in front
+ # of the cursor.
+ deleted = event.current_buffer.delete(count=-event.arg)
+ else:
+ deleted = event.current_buffer.delete_before_cursor(count=event.arg)
+
+ if not deleted:
+ event.app.output.bell()
+
+
+@register("self-insert")
+def self_insert(event: E) -> None:
+ """
+ Insert yourself.
+ """
+ event.current_buffer.insert_text(event.data * event.arg)
+
+
+@register("transpose-chars")
+def transpose_chars(event: E) -> None:
+ """
+ Emulate Emacs transpose-char behavior: at the beginning of the buffer,
+ do nothing. At the end of a line or buffer, swap the characters before
+ the cursor. Otherwise, move the cursor right, and then swap the
+ characters before the cursor.
+ """
+ b = event.current_buffer
+ p = b.cursor_position
+ if p == 0:
+ return
+ elif p == len(b.text) or b.text[p] == "\n":
+ b.swap_characters_before_cursor()
+ else:
+ b.cursor_position += b.document.get_cursor_right_position()
+ b.swap_characters_before_cursor()
+
+
+@register("uppercase-word")
+def uppercase_word(event: E) -> None:
+ """
+ Uppercase the current (or following) word.
+ """
+ buff = event.current_buffer
+
+ for i in range(event.arg):
+ pos = buff.document.find_next_word_ending()
+ words = buff.document.text_after_cursor[:pos]
+ buff.insert_text(words.upper(), overwrite=True)
+
+
+@register("downcase-word")
+def downcase_word(event: E) -> None:
+ """
+ Lowercase the current (or following) word.
+ """
+ buff = event.current_buffer
+
+ for i in range(event.arg): # XXX: not DRY: see meta_c and meta_u!!
+ pos = buff.document.find_next_word_ending()
+ words = buff.document.text_after_cursor[:pos]
+ buff.insert_text(words.lower(), overwrite=True)
+
+
+@register("capitalize-word")
+def capitalize_word(event: E) -> None:
+ """
+ Capitalize the current (or following) word.
+ """
+ buff = event.current_buffer
+
+ for i in range(event.arg):
+ pos = buff.document.find_next_word_ending()
+ words = buff.document.text_after_cursor[:pos]
+ buff.insert_text(words.title(), overwrite=True)
+
+
+@register("quoted-insert")
+def quoted_insert(event: E) -> None:
+ """
+ Add the next character typed to the line verbatim. This is how to insert
+ key sequences like C-q, for example.
+ """
+ event.app.quoted_insert = True
+
+
+#
+# Killing and yanking.
+#
+
+
+@register("kill-line")
+def kill_line(event: E) -> None:
+ """
+ Kill the text from the cursor to the end of the line.
+
+ If we are at the end of the line, this should remove the newline.
+ (That way, it is possible to delete multiple lines by executing this
+ command multiple times.)
+ """
+ buff = event.current_buffer
+ if event.arg < 0:
+ deleted = buff.delete_before_cursor(
+ count=-buff.document.get_start_of_line_position()
+ )
+ else:
+ if buff.document.current_char == "\n":
+ deleted = buff.delete(1)
+ else:
+ deleted = buff.delete(count=buff.document.get_end_of_line_position())
+ event.app.clipboard.set_text(deleted)
+
+
+@register("kill-word")
+def kill_word(event: E) -> None:
+ """
+ Kill from point to the end of the current word, or if between words, to the
+ end of the next word. Word boundaries are the same as forward-word.
+ """
+ buff = event.current_buffer
+ pos = buff.document.find_next_word_ending(count=event.arg)
+
+ if pos:
+ deleted = buff.delete(count=pos)
+
+ if event.is_repeat:
+ deleted = event.app.clipboard.get_data().text + deleted
+
+ event.app.clipboard.set_text(deleted)
+
+
+@register("unix-word-rubout")
+def unix_word_rubout(event: E, WORD: bool = True) -> None:
+ """
+ Kill the word behind point, using whitespace as a word boundary.
+ Usually bound to ControlW.
+ """
+ buff = event.current_buffer
+ pos = buff.document.find_start_of_previous_word(count=event.arg, WORD=WORD)
+
+ if pos is None:
+ # Nothing found? delete until the start of the document. (The
+ # input starts with whitespace and no words were found before the
+ # cursor.)
+ pos = -buff.cursor_position
+
+ if pos:
+ deleted = buff.delete_before_cursor(count=-pos)
+
+ # If the previous key press was also Control-W, concatenate deleted
+ # text.
+ if event.is_repeat:
+ deleted += event.app.clipboard.get_data().text
+
+ event.app.clipboard.set_text(deleted)
+ else:
+ # Nothing to delete. Bell.
+ event.app.output.bell()
+
+
+@register("backward-kill-word")
+def backward_kill_word(event: E) -> None:
+ """
+ Kills the word before point, using "not a letter nor a digit" as a word boundary.
+ Usually bound to M-Del or M-Backspace.
+ """
+ unix_word_rubout(event, WORD=False)
+
+
+@register("delete-horizontal-space")
+def delete_horizontal_space(event: E) -> None:
+ """
+ Delete all spaces and tabs around point.
+ """
+ buff = event.current_buffer
+ text_before_cursor = buff.document.text_before_cursor
+ text_after_cursor = buff.document.text_after_cursor
+
+ delete_before = len(text_before_cursor) - len(text_before_cursor.rstrip("\t "))
+ delete_after = len(text_after_cursor) - len(text_after_cursor.lstrip("\t "))
+
+ buff.delete_before_cursor(count=delete_before)
+ buff.delete(count=delete_after)
+
+
+@register("unix-line-discard")
+def unix_line_discard(event: E) -> None:
+ """
+ Kill backward from the cursor to the beginning of the current line.
+ """
+ buff = event.current_buffer
+
+ if buff.document.cursor_position_col == 0 and buff.document.cursor_position > 0:
+ buff.delete_before_cursor(count=1)
+ else:
+ deleted = buff.delete_before_cursor(
+ count=-buff.document.get_start_of_line_position()
+ )
+ event.app.clipboard.set_text(deleted)
+
+
+@register("yank")
+def yank(event: E) -> None:
+ """
+ Paste before cursor.
+ """
+ event.current_buffer.paste_clipboard_data(
+ event.app.clipboard.get_data(), count=event.arg, paste_mode=PasteMode.EMACS
+ )
+
+
+@register("yank-nth-arg")
+def yank_nth_arg(event: E) -> None:
+ """
+ Insert the first argument of the previous command. With an argument, insert
+ the nth word from the previous command (start counting at 0).
+ """
+ n = event.arg if event.arg_present else None
+ event.current_buffer.yank_nth_arg(n)
+
+
+@register("yank-last-arg")
+def yank_last_arg(event: E) -> None:
+ """
+ Like `yank_nth_arg`, but if no argument has been given, yank the last word
+ of each line.
+ """
+ n = event.arg if event.arg_present else None
+ event.current_buffer.yank_last_arg(n)
+
+
+@register("yank-pop")
+def yank_pop(event: E) -> None:
+ """
+ Rotate the kill ring, and yank the new top. Only works following yank or
+ yank-pop.
+ """
+ buff = event.current_buffer
+ doc_before_paste = buff.document_before_paste
+ clipboard = event.app.clipboard
+
+ if doc_before_paste is not None:
+ buff.document = doc_before_paste
+ clipboard.rotate()
+ buff.paste_clipboard_data(clipboard.get_data(), paste_mode=PasteMode.EMACS)
+
+
+#
+# Completion.
+#
+
+
+@register("complete")
+def complete(event: E) -> None:
+ """
+ Attempt to perform completion.
+ """
+ display_completions_like_readline(event)
+
+
+@register("menu-complete")
+def menu_complete(event: E) -> None:
+ """
+ Generate completions, or go to the next completion. (This is the default
+ way of completing input in prompt_toolkit.)
+ """
+ generate_completions(event)
+
+
+@register("menu-complete-backward")
+def menu_complete_backward(event: E) -> None:
+ """
+ Move backward through the list of possible completions.
+ """
+ event.current_buffer.complete_previous()
+
+
+#
+# Keyboard macros.
+#
+
+
+@register("start-kbd-macro")
+def start_kbd_macro(event: E) -> None:
+ """
+ Begin saving the characters typed into the current keyboard macro.
+ """
+ event.app.emacs_state.start_macro()
+
+
+@register("end-kbd-macro")
+def end_kbd_macro(event: E) -> None:
+ """
+ Stop saving the characters typed into the current keyboard macro and save
+ the definition.
+ """
+ event.app.emacs_state.end_macro()
+
+
+@register("call-last-kbd-macro")
+@key_binding(record_in_macro=False)
+def call_last_kbd_macro(event: E) -> None:
+ """
+ Re-execute the last keyboard macro defined, by making the characters in the
+ macro appear as if typed at the keyboard.
+
+ Notice that we pass `record_in_macro=False`. This ensures that the 'c-x e'
+ key sequence doesn't appear in the recording itself. This function inserts
+ the body of the called macro back into the KeyProcessor, so these keys will
+ be added later on to the macro of their handlers have `record_in_macro=True`.
+ """
+ # Insert the macro.
+ macro = event.app.emacs_state.macro
+
+ if macro:
+ event.app.key_processor.feed_multiple(macro, first=True)
+
+
+@register("print-last-kbd-macro")
+def print_last_kbd_macro(event: E) -> None:
+ """
+ Print the last keyboard macro.
+ """
+
+ # TODO: Make the format suitable for the inputrc file.
+ def print_macro() -> None:
+ macro = event.app.emacs_state.macro
+ if macro:
+ for k in macro:
+ print(k)
+
+ from prompt_toolkit.application.run_in_terminal import run_in_terminal
+
+ run_in_terminal(print_macro)
+
+
+#
+# Miscellaneous Commands.
+#
+
+
+@register("undo")
+def undo(event: E) -> None:
+ """
+ Incremental undo.
+ """
+ event.current_buffer.undo()
+
+
+@register("insert-comment")
+def insert_comment(event: E) -> None:
+ """
+ Without numeric argument, comment all lines.
+ With numeric argument, uncomment all lines.
+ In any case accept the input.
+ """
+ buff = event.current_buffer
+
+ # Transform all lines.
+ if event.arg != 1:
+
+ def change(line: str) -> str:
+ return line[1:] if line.startswith("#") else line
+
+ else:
+
+ def change(line: str) -> str:
+ return "#" + line
+
+ buff.document = Document(
+ text="\n".join(map(change, buff.text.splitlines())), cursor_position=0
+ )
+
+ # Accept input.
+ buff.validate_and_handle()
+
+
+@register("vi-editing-mode")
+def vi_editing_mode(event: E) -> None:
+ """
+ Switch to Vi editing mode.
+ """
+ event.app.editing_mode = EditingMode.VI
+
+
+@register("emacs-editing-mode")
+def emacs_editing_mode(event: E) -> None:
+ """
+ Switch to Emacs editing mode.
+ """
+ event.app.editing_mode = EditingMode.EMACS
+
+
+@register("prefix-meta")
+def prefix_meta(event: E) -> None:
+ """
+ Metafy the next character typed. This is for keyboards without a meta key.
+
+ Sometimes people also want to bind other keys to Meta, e.g. 'jj'::
+
+ key_bindings.add_key_binding('j', 'j', filter=ViInsertMode())(prefix_meta)
+ """
+ # ('first' should be true, because we want to insert it at the current
+ # position in the queue.)
+ event.app.key_processor.feed(KeyPress(Keys.Escape), first=True)
+
+
+@register("operate-and-get-next")
+def operate_and_get_next(event: E) -> None:
+ """
+ Accept the current line for execution and fetch the next line relative to
+ the current line from the history for editing.
+ """
+ buff = event.current_buffer
+ new_index = buff.working_index + 1
+
+ # Accept the current input. (This will also redraw the interface in the
+ # 'done' state.)
+ buff.validate_and_handle()
+
+ # Set the new index at the start of the next run.
+ def set_working_index() -> None:
+ if new_index < len(buff._working_lines):
+ buff.working_index = new_index
+
+ event.app.pre_run_callables.append(set_working_index)
+
+
+@register("edit-and-execute-command")
+def edit_and_execute(event: E) -> None:
+ """
+ Invoke an editor on the current command line, and accept the result.
+ """
+ buff = event.current_buffer
+ buff.open_in_editor(validate_and_handle=True)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/open_in_editor.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/open_in_editor.py
new file mode 100644
index 0000000000000000000000000000000000000000..26b8685b0ae75d7ac05bf4d99be4c44e91a51de1
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/open_in_editor.py
@@ -0,0 +1,52 @@
+"""
+Open in editor key bindings.
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit.filters import emacs_mode, has_selection, vi_navigation_mode
+
+from ..key_bindings import KeyBindings, KeyBindingsBase, merge_key_bindings
+from .named_commands import get_by_name
+
+__all__ = [
+ "load_open_in_editor_bindings",
+ "load_emacs_open_in_editor_bindings",
+ "load_vi_open_in_editor_bindings",
+]
+
+
+def load_open_in_editor_bindings() -> KeyBindingsBase:
+ """
+ Load both the Vi and emacs key bindings for handling edit-and-execute-command.
+ """
+ return merge_key_bindings(
+ [
+ load_emacs_open_in_editor_bindings(),
+ load_vi_open_in_editor_bindings(),
+ ]
+ )
+
+
+def load_emacs_open_in_editor_bindings() -> KeyBindings:
+ """
+ Pressing C-X C-E will open the buffer in an external editor.
+ """
+ key_bindings = KeyBindings()
+
+ key_bindings.add("c-x", "c-e", filter=emacs_mode & ~has_selection)(
+ get_by_name("edit-and-execute-command")
+ )
+
+ return key_bindings
+
+
+def load_vi_open_in_editor_bindings() -> KeyBindings:
+ """
+ Pressing 'v' in navigation mode will open the buffer in an external editor.
+ """
+ key_bindings = KeyBindings()
+ key_bindings.add("v", filter=vi_navigation_mode)(
+ get_by_name("edit-and-execute-command")
+ )
+ return key_bindings
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/page_navigation.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/page_navigation.py
new file mode 100644
index 0000000000000000000000000000000000000000..c490425fe205557108ec8eaac75e9b6e36561a7b
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/page_navigation.py
@@ -0,0 +1,85 @@
+"""
+Key bindings for extra page navigation: bindings for up/down scrolling through
+long pages, like in Emacs or Vi.
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit.filters import buffer_has_focus, emacs_mode, vi_mode
+from prompt_toolkit.key_binding.key_bindings import (
+ ConditionalKeyBindings,
+ KeyBindings,
+ KeyBindingsBase,
+ merge_key_bindings,
+)
+
+from .scroll import (
+ scroll_backward,
+ scroll_forward,
+ scroll_half_page_down,
+ scroll_half_page_up,
+ scroll_one_line_down,
+ scroll_one_line_up,
+ scroll_page_down,
+ scroll_page_up,
+)
+
+__all__ = [
+ "load_page_navigation_bindings",
+ "load_emacs_page_navigation_bindings",
+ "load_vi_page_navigation_bindings",
+]
+
+
+def load_page_navigation_bindings() -> KeyBindingsBase:
+ """
+ Load both the Vi and Emacs bindings for page navigation.
+ """
+ # Only enable when a `Buffer` is focused, otherwise, we would catch keys
+ # when another widget is focused (like for instance `c-d` in a
+ # ptterm.Terminal).
+ return ConditionalKeyBindings(
+ merge_key_bindings(
+ [
+ load_emacs_page_navigation_bindings(),
+ load_vi_page_navigation_bindings(),
+ ]
+ ),
+ buffer_has_focus,
+ )
+
+
+def load_emacs_page_navigation_bindings() -> KeyBindingsBase:
+ """
+ Key bindings, for scrolling up and down through pages.
+ This are separate bindings, because GNU readline doesn't have them.
+ """
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+
+ handle("c-v")(scroll_page_down)
+ handle("pagedown")(scroll_page_down)
+ handle("escape", "v")(scroll_page_up)
+ handle("pageup")(scroll_page_up)
+
+ return ConditionalKeyBindings(key_bindings, emacs_mode)
+
+
+def load_vi_page_navigation_bindings() -> KeyBindingsBase:
+ """
+ Key bindings, for scrolling up and down through pages.
+ This are separate bindings, because GNU readline doesn't have them.
+ """
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+
+ handle("c-f")(scroll_forward)
+ handle("c-b")(scroll_backward)
+ handle("c-d")(scroll_half_page_down)
+ handle("c-u")(scroll_half_page_up)
+ handle("c-e")(scroll_one_line_down)
+ handle("c-y")(scroll_one_line_up)
+ handle("pagedown")(scroll_page_down)
+ handle("pageup")(scroll_page_up)
+
+ return ConditionalKeyBindings(key_bindings, vi_mode)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/scroll.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/scroll.py
new file mode 100644
index 0000000000000000000000000000000000000000..13e44ed4dc947fe7878731fd5063682bac46b281
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/scroll.py
@@ -0,0 +1,190 @@
+"""
+Key bindings, for scrolling up and down through pages.
+
+This are separate bindings, because GNU readline doesn't have them, but
+they are very useful for navigating through long multiline buffers, like in
+Vi, Emacs, etc...
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+
+__all__ = [
+ "scroll_forward",
+ "scroll_backward",
+ "scroll_half_page_up",
+ "scroll_half_page_down",
+ "scroll_one_line_up",
+ "scroll_one_line_down",
+]
+
+E = KeyPressEvent
+
+
+def scroll_forward(event: E, half: bool = False) -> None:
+ """
+ Scroll window down.
+ """
+ w = event.app.layout.current_window
+ b = event.app.current_buffer
+
+ if w and w.render_info:
+ info = w.render_info
+ ui_content = info.ui_content
+
+ # Height to scroll.
+ scroll_height = info.window_height
+ if half:
+ scroll_height //= 2
+
+ # Calculate how many lines is equivalent to that vertical space.
+ y = b.document.cursor_position_row + 1
+ height = 0
+ while y < ui_content.line_count:
+ line_height = info.get_height_for_line(y)
+
+ if height + line_height < scroll_height:
+ height += line_height
+ y += 1
+ else:
+ break
+
+ b.cursor_position = b.document.translate_row_col_to_index(y, 0)
+
+
+def scroll_backward(event: E, half: bool = False) -> None:
+ """
+ Scroll window up.
+ """
+ w = event.app.layout.current_window
+ b = event.app.current_buffer
+
+ if w and w.render_info:
+ info = w.render_info
+
+ # Height to scroll.
+ scroll_height = info.window_height
+ if half:
+ scroll_height //= 2
+
+ # Calculate how many lines is equivalent to that vertical space.
+ y = max(0, b.document.cursor_position_row - 1)
+ height = 0
+ while y > 0:
+ line_height = info.get_height_for_line(y)
+
+ if height + line_height < scroll_height:
+ height += line_height
+ y -= 1
+ else:
+ break
+
+ b.cursor_position = b.document.translate_row_col_to_index(y, 0)
+
+
+def scroll_half_page_down(event: E) -> None:
+ """
+ Same as ControlF, but only scroll half a page.
+ """
+ scroll_forward(event, half=True)
+
+
+def scroll_half_page_up(event: E) -> None:
+ """
+ Same as ControlB, but only scroll half a page.
+ """
+ scroll_backward(event, half=True)
+
+
+def scroll_one_line_down(event: E) -> None:
+ """
+ scroll_offset += 1
+ """
+ w = event.app.layout.current_window
+ b = event.app.current_buffer
+
+ if w:
+ # When the cursor is at the top, move to the next line. (Otherwise, only scroll.)
+ if w.render_info:
+ info = w.render_info
+
+ if w.vertical_scroll < info.content_height - info.window_height:
+ if info.cursor_position.y <= info.configured_scroll_offsets.top:
+ b.cursor_position += b.document.get_cursor_down_position()
+
+ w.vertical_scroll += 1
+
+
+def scroll_one_line_up(event: E) -> None:
+ """
+ scroll_offset -= 1
+ """
+ w = event.app.layout.current_window
+ b = event.app.current_buffer
+
+ if w:
+ # When the cursor is at the bottom, move to the previous line. (Otherwise, only scroll.)
+ if w.render_info:
+ info = w.render_info
+
+ if w.vertical_scroll > 0:
+ first_line_height = info.get_height_for_line(info.first_visible_line())
+
+ cursor_up = info.cursor_position.y - (
+ info.window_height
+ - 1
+ - first_line_height
+ - info.configured_scroll_offsets.bottom
+ )
+
+ # Move cursor up, as many steps as the height of the first line.
+ # TODO: not entirely correct yet, in case of line wrapping and many long lines.
+ for _ in range(max(0, cursor_up)):
+ b.cursor_position += b.document.get_cursor_up_position()
+
+ # Scroll window
+ w.vertical_scroll -= 1
+
+
+def scroll_page_down(event: E) -> None:
+ """
+ Scroll page down. (Prefer the cursor at the top of the page, after scrolling.)
+ """
+ w = event.app.layout.current_window
+ b = event.app.current_buffer
+
+ if w and w.render_info:
+ # Scroll down one page.
+ line_index = max(w.render_info.last_visible_line(), w.vertical_scroll + 1)
+ w.vertical_scroll = line_index
+
+ b.cursor_position = b.document.translate_row_col_to_index(line_index, 0)
+ b.cursor_position += b.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+
+
+def scroll_page_up(event: E) -> None:
+ """
+ Scroll page up. (Prefer the cursor at the bottom of the page, after scrolling.)
+ """
+ w = event.app.layout.current_window
+ b = event.app.current_buffer
+
+ if w and w.render_info:
+ # Put cursor at the first visible line. (But make sure that the cursor
+ # moves at least one line up.)
+ line_index = max(
+ 0,
+ min(w.render_info.first_visible_line(), b.document.cursor_position_row - 1),
+ )
+
+ b.cursor_position = b.document.translate_row_col_to_index(line_index, 0)
+ b.cursor_position += b.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+
+ # Set the scroll offset. We can safely set it to zero; the Window will
+ # make sure that it scrolls at least until the cursor becomes visible.
+ w.vertical_scroll = 0
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/search.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/search.py
new file mode 100644
index 0000000000000000000000000000000000000000..a57c52e4f237baa5a4c54de2d5b70447b3aaceb0
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/search.py
@@ -0,0 +1,96 @@
+"""
+Search related key bindings.
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit import search
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.filters import Condition, control_is_searchable, is_searching
+from prompt_toolkit.key_binding.key_processor import KeyPressEvent
+
+from ..key_bindings import key_binding
+
+__all__ = [
+ "abort_search",
+ "accept_search",
+ "start_reverse_incremental_search",
+ "start_forward_incremental_search",
+ "reverse_incremental_search",
+ "forward_incremental_search",
+ "accept_search_and_accept_input",
+]
+
+E = KeyPressEvent
+
+
+@key_binding(filter=is_searching)
+def abort_search(event: E) -> None:
+ """
+ Abort an incremental search and restore the original
+ line.
+ (Usually bound to ControlG/ControlC.)
+ """
+ search.stop_search()
+
+
+@key_binding(filter=is_searching)
+def accept_search(event: E) -> None:
+ """
+ When enter pressed in isearch, quit isearch mode. (Multiline
+ isearch would be too complicated.)
+ (Usually bound to Enter.)
+ """
+ search.accept_search()
+
+
+@key_binding(filter=control_is_searchable)
+def start_reverse_incremental_search(event: E) -> None:
+ """
+ Enter reverse incremental search.
+ (Usually ControlR.)
+ """
+ search.start_search(direction=search.SearchDirection.BACKWARD)
+
+
+@key_binding(filter=control_is_searchable)
+def start_forward_incremental_search(event: E) -> None:
+ """
+ Enter forward incremental search.
+ (Usually ControlS.)
+ """
+ search.start_search(direction=search.SearchDirection.FORWARD)
+
+
+@key_binding(filter=is_searching)
+def reverse_incremental_search(event: E) -> None:
+ """
+ Apply reverse incremental search, but keep search buffer focused.
+ """
+ search.do_incremental_search(search.SearchDirection.BACKWARD, count=event.arg)
+
+
+@key_binding(filter=is_searching)
+def forward_incremental_search(event: E) -> None:
+ """
+ Apply forward incremental search, but keep search buffer focused.
+ """
+ search.do_incremental_search(search.SearchDirection.FORWARD, count=event.arg)
+
+
+@Condition
+def _previous_buffer_is_returnable() -> bool:
+ """
+ True if the previously focused buffer has a return handler.
+ """
+ prev_control = get_app().layout.search_target_buffer_control
+ return bool(prev_control and prev_control.buffer.is_returnable)
+
+
+@key_binding(filter=is_searching & _previous_buffer_is_returnable)
+def accept_search_and_accept_input(event: E) -> None:
+ """
+ Accept the search operation first, then accept the input.
+ """
+ search.accept_search()
+ event.current_buffer.validate_and_handle()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/vi.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/vi.py
new file mode 100644
index 0000000000000000000000000000000000000000..d68a31f29a0635c107096b28271a8e68dcbaf894
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/bindings/vi.py
@@ -0,0 +1,2233 @@
+# pylint: disable=function-redefined
+from __future__ import annotations
+
+import codecs
+import string
+from enum import Enum
+from itertools import accumulate
+from typing import Callable, Iterable, Tuple, TypeVar
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.buffer import Buffer, indent, reshape_text, unindent
+from prompt_toolkit.clipboard import ClipboardData
+from prompt_toolkit.document import Document
+from prompt_toolkit.filters import (
+ Always,
+ Condition,
+ Filter,
+ has_arg,
+ is_read_only,
+ is_searching,
+)
+from prompt_toolkit.filters.app import (
+ in_paste_mode,
+ is_multiline,
+ vi_digraph_mode,
+ vi_insert_mode,
+ vi_insert_multiple_mode,
+ vi_mode,
+ vi_navigation_mode,
+ vi_recording_macro,
+ vi_replace_mode,
+ vi_replace_single_mode,
+ vi_search_direction_reversed,
+ vi_selection_mode,
+ vi_waiting_for_text_object_mode,
+)
+from prompt_toolkit.input.vt100_parser import Vt100Parser
+from prompt_toolkit.key_binding.digraphs import DIGRAPHS
+from prompt_toolkit.key_binding.key_processor import KeyPress, KeyPressEvent
+from prompt_toolkit.key_binding.vi_state import CharacterFind, InputMode
+from prompt_toolkit.keys import Keys
+from prompt_toolkit.search import SearchDirection
+from prompt_toolkit.selection import PasteMode, SelectionState, SelectionType
+
+from ..key_bindings import ConditionalKeyBindings, KeyBindings, KeyBindingsBase
+from .named_commands import get_by_name
+
+__all__ = [
+ "load_vi_bindings",
+ "load_vi_search_bindings",
+]
+
+E = KeyPressEvent
+
+ascii_lowercase = string.ascii_lowercase
+
+vi_register_names = ascii_lowercase + "0123456789"
+
+
+class TextObjectType(Enum):
+ EXCLUSIVE = "EXCLUSIVE"
+ INCLUSIVE = "INCLUSIVE"
+ LINEWISE = "LINEWISE"
+ BLOCK = "BLOCK"
+
+
+class TextObject:
+ """
+ Return struct for functions wrapped in ``text_object``.
+ Both `start` and `end` are relative to the current cursor position.
+ """
+
+ def __init__(
+ self, start: int, end: int = 0, type: TextObjectType = TextObjectType.EXCLUSIVE
+ ):
+ self.start = start
+ self.end = end
+ self.type = type
+
+ @property
+ def selection_type(self) -> SelectionType:
+ if self.type == TextObjectType.LINEWISE:
+ return SelectionType.LINES
+ if self.type == TextObjectType.BLOCK:
+ return SelectionType.BLOCK
+ else:
+ return SelectionType.CHARACTERS
+
+ def sorted(self) -> tuple[int, int]:
+ """
+ Return a (start, end) tuple where start <= end.
+ """
+ if self.start < self.end:
+ return self.start, self.end
+ else:
+ return self.end, self.start
+
+ def operator_range(self, document: Document) -> tuple[int, int]:
+ """
+ Return a (start, end) tuple with start <= end that indicates the range
+ operators should operate on.
+ `buffer` is used to get start and end of line positions.
+
+ This should return something that can be used in a slice, so the `end`
+ position is *not* included.
+ """
+ start, end = self.sorted()
+ doc = document
+
+ if (
+ self.type == TextObjectType.EXCLUSIVE
+ and doc.translate_index_to_position(end + doc.cursor_position)[1] == 0
+ ):
+ # If the motion is exclusive and the end of motion is on the first
+ # column, the end position becomes end of previous line.
+ end -= 1
+ if self.type == TextObjectType.INCLUSIVE:
+ end += 1
+ if self.type == TextObjectType.LINEWISE:
+ # Select whole lines
+ row, col = doc.translate_index_to_position(start + doc.cursor_position)
+ start = doc.translate_row_col_to_index(row, 0) - doc.cursor_position
+ row, col = doc.translate_index_to_position(end + doc.cursor_position)
+ end = (
+ doc.translate_row_col_to_index(row, len(doc.lines[row]))
+ - doc.cursor_position
+ )
+ return start, end
+
+ def get_line_numbers(self, buffer: Buffer) -> tuple[int, int]:
+ """
+ Return a (start_line, end_line) pair.
+ """
+ # Get absolute cursor positions from the text object.
+ from_, to = self.operator_range(buffer.document)
+ from_ += buffer.cursor_position
+ to += buffer.cursor_position
+
+ # Take the start of the lines.
+ from_, _ = buffer.document.translate_index_to_position(from_)
+ to, _ = buffer.document.translate_index_to_position(to)
+
+ return from_, to
+
+ def cut(self, buffer: Buffer) -> tuple[Document, ClipboardData]:
+ """
+ Turn text object into `ClipboardData` instance.
+ """
+ from_, to = self.operator_range(buffer.document)
+
+ from_ += buffer.cursor_position
+ to += buffer.cursor_position
+
+ # For Vi mode, the SelectionState does include the upper position,
+ # while `self.operator_range` does not. So, go one to the left, unless
+ # we're in the line mode, then we don't want to risk going to the
+ # previous line, and missing one line in the selection.
+ if self.type != TextObjectType.LINEWISE:
+ to -= 1
+
+ document = Document(
+ buffer.text,
+ to,
+ SelectionState(original_cursor_position=from_, type=self.selection_type),
+ )
+
+ new_document, clipboard_data = document.cut_selection()
+ return new_document, clipboard_data
+
+
+# Typevar for any text object function:
+TextObjectFunction = Callable[[E], TextObject]
+_TOF = TypeVar("_TOF", bound=TextObjectFunction)
+
+
+def create_text_object_decorator(
+ key_bindings: KeyBindings,
+) -> Callable[..., Callable[[_TOF], _TOF]]:
+ """
+ Create a decorator that can be used to register Vi text object implementations.
+ """
+
+ def text_object_decorator(
+ *keys: Keys | str,
+ filter: Filter = Always(),
+ no_move_handler: bool = False,
+ no_selection_handler: bool = False,
+ eager: bool = False,
+ ) -> Callable[[_TOF], _TOF]:
+ """
+ Register a text object function.
+
+ Usage::
+
+ @text_object('w', filter=..., no_move_handler=False)
+ def handler(event):
+ # Return a text object for this key.
+ return TextObject(...)
+
+ :param no_move_handler: Disable the move handler in navigation mode.
+ (It's still active in selection mode.)
+ """
+
+ def decorator(text_object_func: _TOF) -> _TOF:
+ @key_bindings.add(
+ *keys, filter=vi_waiting_for_text_object_mode & filter, eager=eager
+ )
+ def _apply_operator_to_text_object(event: E) -> None:
+ # Arguments are multiplied.
+ vi_state = event.app.vi_state
+ event._arg = str((vi_state.operator_arg or 1) * (event.arg or 1))
+
+ # Call the text object handler.
+ text_obj = text_object_func(event)
+
+ # Get the operator function.
+ # (Should never be None here, given the
+ # `vi_waiting_for_text_object_mode` filter state.)
+ operator_func = vi_state.operator_func
+
+ if text_obj is not None and operator_func is not None:
+ # Call the operator function with the text object.
+ operator_func(event, text_obj)
+
+ # Clear operator.
+ event.app.vi_state.operator_func = None
+ event.app.vi_state.operator_arg = None
+
+ # Register a move operation. (Doesn't need an operator.)
+ if not no_move_handler:
+
+ @key_bindings.add(
+ *keys,
+ filter=~vi_waiting_for_text_object_mode
+ & filter
+ & vi_navigation_mode,
+ eager=eager,
+ )
+ def _move_in_navigation_mode(event: E) -> None:
+ """
+ Move handler for navigation mode.
+ """
+ text_object = text_object_func(event)
+ event.current_buffer.cursor_position += text_object.start
+
+ # Register a move selection operation.
+ if not no_selection_handler:
+
+ @key_bindings.add(
+ *keys,
+ filter=~vi_waiting_for_text_object_mode
+ & filter
+ & vi_selection_mode,
+ eager=eager,
+ )
+ def _move_in_selection_mode(event: E) -> None:
+ """
+ Move handler for selection mode.
+ """
+ text_object = text_object_func(event)
+ buff = event.current_buffer
+ selection_state = buff.selection_state
+
+ if selection_state is None:
+ return # Should not happen, because of the `vi_selection_mode` filter.
+
+ # When the text object has both a start and end position, like 'i(' or 'iw',
+ # Turn this into a selection, otherwise the cursor.
+ if text_object.end:
+ # Take selection positions from text object.
+ start, end = text_object.operator_range(buff.document)
+ start += buff.cursor_position
+ end += buff.cursor_position
+
+ selection_state.original_cursor_position = start
+ buff.cursor_position = end
+
+ # Take selection type from text object.
+ if text_object.type == TextObjectType.LINEWISE:
+ selection_state.type = SelectionType.LINES
+ else:
+ selection_state.type = SelectionType.CHARACTERS
+ else:
+ event.current_buffer.cursor_position += text_object.start
+
+ # Make it possible to chain @text_object decorators.
+ return text_object_func
+
+ return decorator
+
+ return text_object_decorator
+
+
+# Typevar for any operator function:
+OperatorFunction = Callable[[E, TextObject], None]
+_OF = TypeVar("_OF", bound=OperatorFunction)
+
+
+def create_operator_decorator(
+ key_bindings: KeyBindings,
+) -> Callable[..., Callable[[_OF], _OF]]:
+ """
+ Create a decorator that can be used for registering Vi operators.
+ """
+
+ def operator_decorator(
+ *keys: Keys | str, filter: Filter = Always(), eager: bool = False
+ ) -> Callable[[_OF], _OF]:
+ """
+ Register a Vi operator.
+
+ Usage::
+
+ @operator('d', filter=...)
+ def handler(event, text_object):
+ # Do something with the text object here.
+ """
+
+ def decorator(operator_func: _OF) -> _OF:
+ @key_bindings.add(
+ *keys,
+ filter=~vi_waiting_for_text_object_mode & filter & vi_navigation_mode,
+ eager=eager,
+ )
+ def _operator_in_navigation(event: E) -> None:
+ """
+ Handle operator in navigation mode.
+ """
+ # When this key binding is matched, only set the operator
+ # function in the ViState. We should execute it after a text
+ # object has been received.
+ event.app.vi_state.operator_func = operator_func
+ event.app.vi_state.operator_arg = event.arg
+
+ @key_bindings.add(
+ *keys,
+ filter=~vi_waiting_for_text_object_mode & filter & vi_selection_mode,
+ eager=eager,
+ )
+ def _operator_in_selection(event: E) -> None:
+ """
+ Handle operator in selection mode.
+ """
+ buff = event.current_buffer
+ selection_state = buff.selection_state
+
+ if selection_state is not None:
+ # Create text object from selection.
+ if selection_state.type == SelectionType.LINES:
+ text_obj_type = TextObjectType.LINEWISE
+ elif selection_state.type == SelectionType.BLOCK:
+ text_obj_type = TextObjectType.BLOCK
+ else:
+ text_obj_type = TextObjectType.INCLUSIVE
+
+ text_object = TextObject(
+ selection_state.original_cursor_position - buff.cursor_position,
+ type=text_obj_type,
+ )
+
+ # Execute operator.
+ operator_func(event, text_object)
+
+ # Quit selection mode.
+ buff.selection_state = None
+
+ return operator_func
+
+ return decorator
+
+ return operator_decorator
+
+
+@Condition
+def is_returnable() -> bool:
+ return get_app().current_buffer.is_returnable
+
+
+@Condition
+def in_block_selection() -> bool:
+ buff = get_app().current_buffer
+ return bool(
+ buff.selection_state and buff.selection_state.type == SelectionType.BLOCK
+ )
+
+
+@Condition
+def digraph_symbol_1_given() -> bool:
+ return get_app().vi_state.digraph_symbol1 is not None
+
+
+@Condition
+def search_buffer_is_empty() -> bool:
+ "Returns True when the search buffer is empty."
+ return get_app().current_buffer.text == ""
+
+
+@Condition
+def tilde_operator() -> bool:
+ return get_app().vi_state.tilde_operator
+
+
+def load_vi_bindings() -> KeyBindingsBase:
+ """
+ Vi extensions.
+
+ # Overview of Readline Vi commands:
+ # http://www.catonmat.net/download/bash-vi-editing-mode-cheat-sheet.pdf
+ """
+ # Note: Some key bindings have the "~IsReadOnly()" filter added. This
+ # prevents the handler to be executed when the focus is on a
+ # read-only buffer.
+ # This is however only required for those that change the ViState to
+ # INSERT mode. The `Buffer` class itself throws the
+ # `EditReadOnlyBuffer` exception for any text operations which is
+ # handled correctly. There is no need to add "~IsReadOnly" to all key
+ # bindings that do text manipulation.
+
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+
+ # (Note: Always take the navigation bindings in read-only mode, even when
+ # ViState says different.)
+
+ TransformFunction = Tuple[Tuple[str, ...], Filter, Callable[[str], str]]
+
+ vi_transform_functions: list[TransformFunction] = [
+ # Rot 13 transformation
+ (
+ ("g", "?"),
+ Always(),
+ lambda string: codecs.encode(string, "rot_13"),
+ ),
+ # To lowercase
+ (("g", "u"), Always(), lambda string: string.lower()),
+ # To uppercase.
+ (("g", "U"), Always(), lambda string: string.upper()),
+ # Swap case.
+ (("g", "~"), Always(), lambda string: string.swapcase()),
+ (
+ ("~",),
+ tilde_operator,
+ lambda string: string.swapcase(),
+ ),
+ ]
+
+ # Insert a character literally (quoted insert).
+ handle("c-v", filter=vi_insert_mode)(get_by_name("quoted-insert"))
+
+ @handle("escape")
+ def _back_to_navigation(event: E) -> None:
+ """
+ Escape goes to vi navigation mode.
+ """
+ buffer = event.current_buffer
+ vi_state = event.app.vi_state
+
+ if vi_state.input_mode in (InputMode.INSERT, InputMode.REPLACE):
+ buffer.cursor_position += buffer.document.get_cursor_left_position()
+
+ vi_state.input_mode = InputMode.NAVIGATION
+
+ if bool(buffer.selection_state):
+ buffer.exit_selection()
+
+ @handle("k", filter=vi_selection_mode)
+ def _up_in_selection(event: E) -> None:
+ """
+ Arrow up in selection mode.
+ """
+ event.current_buffer.cursor_up(count=event.arg)
+
+ @handle("j", filter=vi_selection_mode)
+ def _down_in_selection(event: E) -> None:
+ """
+ Arrow down in selection mode.
+ """
+ event.current_buffer.cursor_down(count=event.arg)
+
+ @handle("up", filter=vi_navigation_mode)
+ @handle("c-p", filter=vi_navigation_mode)
+ def _up_in_navigation(event: E) -> None:
+ """
+ Arrow up and ControlP in navigation mode go up.
+ """
+ event.current_buffer.auto_up(count=event.arg)
+
+ @handle("k", filter=vi_navigation_mode)
+ def _go_up(event: E) -> None:
+ """
+ Go up, but if we enter a new history entry, move to the start of the
+ line.
+ """
+ event.current_buffer.auto_up(
+ count=event.arg, go_to_start_of_line_if_history_changes=True
+ )
+
+ @handle("down", filter=vi_navigation_mode)
+ @handle("c-n", filter=vi_navigation_mode)
+ def _go_down(event: E) -> None:
+ """
+ Arrow down and Control-N in navigation mode.
+ """
+ event.current_buffer.auto_down(count=event.arg)
+
+ @handle("j", filter=vi_navigation_mode)
+ def _go_down2(event: E) -> None:
+ """
+ Go down, but if we enter a new history entry, go to the start of the line.
+ """
+ event.current_buffer.auto_down(
+ count=event.arg, go_to_start_of_line_if_history_changes=True
+ )
+
+ @handle("backspace", filter=vi_navigation_mode)
+ def _go_left(event: E) -> None:
+ """
+ In navigation-mode, move cursor.
+ """
+ event.current_buffer.cursor_position += (
+ event.current_buffer.document.get_cursor_left_position(count=event.arg)
+ )
+
+ @handle("c-n", filter=vi_insert_mode)
+ def _complete_next(event: E) -> None:
+ b = event.current_buffer
+
+ if b.complete_state:
+ b.complete_next()
+ else:
+ b.start_completion(select_first=True)
+
+ @handle("c-p", filter=vi_insert_mode)
+ def _complete_prev(event: E) -> None:
+ """
+ Control-P: To previous completion.
+ """
+ b = event.current_buffer
+
+ if b.complete_state:
+ b.complete_previous()
+ else:
+ b.start_completion(select_last=True)
+
+ @handle("c-g", filter=vi_insert_mode)
+ @handle("c-y", filter=vi_insert_mode)
+ def _accept_completion(event: E) -> None:
+ """
+ Accept current completion.
+ """
+ event.current_buffer.complete_state = None
+
+ @handle("c-e", filter=vi_insert_mode)
+ def _cancel_completion(event: E) -> None:
+ """
+ Cancel completion. Go back to originally typed text.
+ """
+ event.current_buffer.cancel_completion()
+
+ # In navigation mode, pressing enter will always return the input.
+ handle("enter", filter=vi_navigation_mode & is_returnable)(
+ get_by_name("accept-line")
+ )
+
+ # In insert mode, also accept input when enter is pressed, and the buffer
+ # has been marked as single line.
+ handle("enter", filter=is_returnable & ~is_multiline)(get_by_name("accept-line"))
+
+ @handle("enter", filter=~is_returnable & vi_navigation_mode)
+ def _start_of_next_line(event: E) -> None:
+ """
+ Go to the beginning of next line.
+ """
+ b = event.current_buffer
+ b.cursor_down(count=event.arg)
+ b.cursor_position += b.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+
+ # ** In navigation mode **
+
+ # List of navigation commands: http://hea-www.harvard.edu/~fine/Tech/vi.html
+
+ @handle("insert", filter=vi_navigation_mode)
+ def _insert_mode(event: E) -> None:
+ """
+ Pressing the Insert key.
+ """
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("insert", filter=vi_insert_mode)
+ def _navigation_mode(event: E) -> None:
+ """
+ Pressing the Insert key.
+ """
+ event.app.vi_state.input_mode = InputMode.NAVIGATION
+
+ @handle("a", filter=vi_navigation_mode & ~is_read_only)
+ # ~IsReadOnly, because we want to stay in navigation mode for
+ # read-only buffers.
+ def _a(event: E) -> None:
+ event.current_buffer.cursor_position += (
+ event.current_buffer.document.get_cursor_right_position()
+ )
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("A", filter=vi_navigation_mode & ~is_read_only)
+ def _A(event: E) -> None:
+ event.current_buffer.cursor_position += (
+ event.current_buffer.document.get_end_of_line_position()
+ )
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("C", filter=vi_navigation_mode & ~is_read_only)
+ def _change_until_end_of_line(event: E) -> None:
+ """
+ Change to end of line.
+ Same as 'c$' (which is implemented elsewhere.)
+ """
+ buffer = event.current_buffer
+
+ deleted = buffer.delete(count=buffer.document.get_end_of_line_position())
+ event.app.clipboard.set_text(deleted)
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("c", "c", filter=vi_navigation_mode & ~is_read_only)
+ @handle("S", filter=vi_navigation_mode & ~is_read_only)
+ def _change_current_line(event: E) -> None: # TODO: implement 'arg'
+ """
+ Change current line
+ """
+ buffer = event.current_buffer
+
+ # We copy the whole line.
+ data = ClipboardData(buffer.document.current_line, SelectionType.LINES)
+ event.app.clipboard.set_data(data)
+
+ # But we delete after the whitespace
+ buffer.cursor_position += buffer.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+ buffer.delete(count=buffer.document.get_end_of_line_position())
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("D", filter=vi_navigation_mode)
+ def _delete_until_end_of_line(event: E) -> None:
+ """
+ Delete from cursor position until the end of the line.
+ """
+ buffer = event.current_buffer
+ deleted = buffer.delete(count=buffer.document.get_end_of_line_position())
+ event.app.clipboard.set_text(deleted)
+
+ @handle("d", "d", filter=vi_navigation_mode)
+ def _delete_line(event: E) -> None:
+ """
+ Delete line. (Or the following 'n' lines.)
+ """
+ buffer = event.current_buffer
+
+ # Split string in before/deleted/after text.
+ lines = buffer.document.lines
+
+ before = "\n".join(lines[: buffer.document.cursor_position_row])
+ deleted = "\n".join(
+ lines[
+ buffer.document.cursor_position_row : buffer.document.cursor_position_row
+ + event.arg
+ ]
+ )
+ after = "\n".join(lines[buffer.document.cursor_position_row + event.arg :])
+
+ # Set new text.
+ if before and after:
+ before = before + "\n"
+
+ # Set text and cursor position.
+ buffer.document = Document(
+ text=before + after,
+ # Cursor At the start of the first 'after' line, after the leading whitespace.
+ cursor_position=len(before) + len(after) - len(after.lstrip(" ")),
+ )
+
+ # Set clipboard data
+ event.app.clipboard.set_data(ClipboardData(deleted, SelectionType.LINES))
+
+ @handle("x", filter=vi_selection_mode)
+ def _cut(event: E) -> None:
+ """
+ Cut selection.
+ ('x' is not an operator.)
+ """
+ clipboard_data = event.current_buffer.cut_selection()
+ event.app.clipboard.set_data(clipboard_data)
+
+ @handle("i", filter=vi_navigation_mode & ~is_read_only)
+ def _i(event: E) -> None:
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("I", filter=vi_navigation_mode & ~is_read_only)
+ def _I(event: E) -> None:
+ event.app.vi_state.input_mode = InputMode.INSERT
+ event.current_buffer.cursor_position += (
+ event.current_buffer.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+ )
+
+ @handle("I", filter=in_block_selection & ~is_read_only)
+ def insert_in_block_selection(event: E, after: bool = False) -> None:
+ """
+ Insert in block selection mode.
+ """
+ buff = event.current_buffer
+
+ # Store all cursor positions.
+ positions = []
+
+ if after:
+
+ def get_pos(from_to: tuple[int, int]) -> int:
+ return from_to[1]
+
+ else:
+
+ def get_pos(from_to: tuple[int, int]) -> int:
+ return from_to[0]
+
+ for i, from_to in enumerate(buff.document.selection_ranges()):
+ positions.append(get_pos(from_to))
+ if i == 0:
+ buff.cursor_position = get_pos(from_to)
+
+ buff.multiple_cursor_positions = positions
+
+ # Go to 'INSERT_MULTIPLE' mode.
+ event.app.vi_state.input_mode = InputMode.INSERT_MULTIPLE
+ buff.exit_selection()
+
+ @handle("A", filter=in_block_selection & ~is_read_only)
+ def _append_after_block(event: E) -> None:
+ insert_in_block_selection(event, after=True)
+
+ @handle("J", filter=vi_navigation_mode & ~is_read_only)
+ def _join(event: E) -> None:
+ """
+ Join lines.
+ """
+ for i in range(event.arg):
+ event.current_buffer.join_next_line()
+
+ @handle("g", "J", filter=vi_navigation_mode & ~is_read_only)
+ def _join_nospace(event: E) -> None:
+ """
+ Join lines without space.
+ """
+ for i in range(event.arg):
+ event.current_buffer.join_next_line(separator="")
+
+ @handle("J", filter=vi_selection_mode & ~is_read_only)
+ def _join_selection(event: E) -> None:
+ """
+ Join selected lines.
+ """
+ event.current_buffer.join_selected_lines()
+
+ @handle("g", "J", filter=vi_selection_mode & ~is_read_only)
+ def _join_selection_nospace(event: E) -> None:
+ """
+ Join selected lines without space.
+ """
+ event.current_buffer.join_selected_lines(separator="")
+
+ @handle("p", filter=vi_navigation_mode)
+ def _paste(event: E) -> None:
+ """
+ Paste after
+ """
+ event.current_buffer.paste_clipboard_data(
+ event.app.clipboard.get_data(),
+ count=event.arg,
+ paste_mode=PasteMode.VI_AFTER,
+ )
+
+ @handle("P", filter=vi_navigation_mode)
+ def _paste_before(event: E) -> None:
+ """
+ Paste before
+ """
+ event.current_buffer.paste_clipboard_data(
+ event.app.clipboard.get_data(),
+ count=event.arg,
+ paste_mode=PasteMode.VI_BEFORE,
+ )
+
+ @handle('"', Keys.Any, "p", filter=vi_navigation_mode)
+ def _paste_register(event: E) -> None:
+ """
+ Paste from named register.
+ """
+ c = event.key_sequence[1].data
+ if c in vi_register_names:
+ data = event.app.vi_state.named_registers.get(c)
+ if data:
+ event.current_buffer.paste_clipboard_data(
+ data, count=event.arg, paste_mode=PasteMode.VI_AFTER
+ )
+
+ @handle('"', Keys.Any, "P", filter=vi_navigation_mode)
+ def _paste_register_before(event: E) -> None:
+ """
+ Paste (before) from named register.
+ """
+ c = event.key_sequence[1].data
+ if c in vi_register_names:
+ data = event.app.vi_state.named_registers.get(c)
+ if data:
+ event.current_buffer.paste_clipboard_data(
+ data, count=event.arg, paste_mode=PasteMode.VI_BEFORE
+ )
+
+ @handle("r", filter=vi_navigation_mode)
+ def _replace(event: E) -> None:
+ """
+ Go to 'replace-single'-mode.
+ """
+ event.app.vi_state.input_mode = InputMode.REPLACE_SINGLE
+
+ @handle("R", filter=vi_navigation_mode)
+ def _replace_mode(event: E) -> None:
+ """
+ Go to 'replace'-mode.
+ """
+ event.app.vi_state.input_mode = InputMode.REPLACE
+
+ @handle("s", filter=vi_navigation_mode & ~is_read_only)
+ def _substitute(event: E) -> None:
+ """
+ Substitute with new text
+ (Delete character(s) and go to insert mode.)
+ """
+ text = event.current_buffer.delete(count=event.arg)
+ event.app.clipboard.set_text(text)
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("u", filter=vi_navigation_mode, save_before=(lambda e: False))
+ def _undo(event: E) -> None:
+ for i in range(event.arg):
+ event.current_buffer.undo()
+
+ @handle("V", filter=vi_navigation_mode)
+ def _visual_line(event: E) -> None:
+ """
+ Start lines selection.
+ """
+ event.current_buffer.start_selection(selection_type=SelectionType.LINES)
+
+ @handle("c-v", filter=vi_navigation_mode)
+ def _visual_block(event: E) -> None:
+ """
+ Enter block selection mode.
+ """
+ event.current_buffer.start_selection(selection_type=SelectionType.BLOCK)
+
+ @handle("V", filter=vi_selection_mode)
+ def _visual_line2(event: E) -> None:
+ """
+ Exit line selection mode, or go from non line selection mode to line
+ selection mode.
+ """
+ selection_state = event.current_buffer.selection_state
+
+ if selection_state is not None:
+ if selection_state.type != SelectionType.LINES:
+ selection_state.type = SelectionType.LINES
+ else:
+ event.current_buffer.exit_selection()
+
+ @handle("v", filter=vi_navigation_mode)
+ def _visual(event: E) -> None:
+ """
+ Enter character selection mode.
+ """
+ event.current_buffer.start_selection(selection_type=SelectionType.CHARACTERS)
+
+ @handle("v", filter=vi_selection_mode)
+ def _visual2(event: E) -> None:
+ """
+ Exit character selection mode, or go from non-character-selection mode
+ to character selection mode.
+ """
+ selection_state = event.current_buffer.selection_state
+
+ if selection_state is not None:
+ if selection_state.type != SelectionType.CHARACTERS:
+ selection_state.type = SelectionType.CHARACTERS
+ else:
+ event.current_buffer.exit_selection()
+
+ @handle("c-v", filter=vi_selection_mode)
+ def _visual_block2(event: E) -> None:
+ """
+ Exit block selection mode, or go from non block selection mode to block
+ selection mode.
+ """
+ selection_state = event.current_buffer.selection_state
+
+ if selection_state is not None:
+ if selection_state.type != SelectionType.BLOCK:
+ selection_state.type = SelectionType.BLOCK
+ else:
+ event.current_buffer.exit_selection()
+
+ @handle("a", "w", filter=vi_selection_mode)
+ @handle("a", "W", filter=vi_selection_mode)
+ def _visual_auto_word(event: E) -> None:
+ """
+ Switch from visual linewise mode to visual characterwise mode.
+ """
+ buffer = event.current_buffer
+
+ if (
+ buffer.selection_state
+ and buffer.selection_state.type == SelectionType.LINES
+ ):
+ buffer.selection_state.type = SelectionType.CHARACTERS
+
+ @handle("x", filter=vi_navigation_mode)
+ def _delete(event: E) -> None:
+ """
+ Delete character.
+ """
+ buff = event.current_buffer
+ count = min(event.arg, len(buff.document.current_line_after_cursor))
+ if count:
+ text = event.current_buffer.delete(count=count)
+ event.app.clipboard.set_text(text)
+
+ @handle("X", filter=vi_navigation_mode)
+ def _delete_before_cursor(event: E) -> None:
+ buff = event.current_buffer
+ count = min(event.arg, len(buff.document.current_line_before_cursor))
+ if count:
+ text = event.current_buffer.delete_before_cursor(count=count)
+ event.app.clipboard.set_text(text)
+
+ @handle("y", "y", filter=vi_navigation_mode)
+ @handle("Y", filter=vi_navigation_mode)
+ def _yank_line(event: E) -> None:
+ """
+ Yank the whole line.
+ """
+ text = "\n".join(event.current_buffer.document.lines_from_current[: event.arg])
+ event.app.clipboard.set_data(ClipboardData(text, SelectionType.LINES))
+
+ @handle("+", filter=vi_navigation_mode)
+ def _next_line(event: E) -> None:
+ """
+ Move to first non whitespace of next line
+ """
+ buffer = event.current_buffer
+ buffer.cursor_position += buffer.document.get_cursor_down_position(
+ count=event.arg
+ )
+ buffer.cursor_position += buffer.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+
+ @handle("-", filter=vi_navigation_mode)
+ def _prev_line(event: E) -> None:
+ """
+ Move to first non whitespace of previous line
+ """
+ buffer = event.current_buffer
+ buffer.cursor_position += buffer.document.get_cursor_up_position(
+ count=event.arg
+ )
+ buffer.cursor_position += buffer.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+
+ @handle(">", ">", filter=vi_navigation_mode)
+ @handle("c-t", filter=vi_insert_mode)
+ def _indent(event: E) -> None:
+ """
+ Indent lines.
+ """
+ buffer = event.current_buffer
+ current_row = buffer.document.cursor_position_row
+ indent(buffer, current_row, current_row + event.arg)
+
+ @handle("<", "<", filter=vi_navigation_mode)
+ @handle("c-d", filter=vi_insert_mode)
+ def _unindent(event: E) -> None:
+ """
+ Unindent lines.
+ """
+ current_row = event.current_buffer.document.cursor_position_row
+ unindent(event.current_buffer, current_row, current_row + event.arg)
+
+ @handle("O", filter=vi_navigation_mode & ~is_read_only)
+ def _open_above(event: E) -> None:
+ """
+ Open line above and enter insertion mode
+ """
+ event.current_buffer.insert_line_above(copy_margin=not in_paste_mode())
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("o", filter=vi_navigation_mode & ~is_read_only)
+ def _open_below(event: E) -> None:
+ """
+ Open line below and enter insertion mode
+ """
+ event.current_buffer.insert_line_below(copy_margin=not in_paste_mode())
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ @handle("~", filter=vi_navigation_mode)
+ def _reverse_case(event: E) -> None:
+ """
+ Reverse case of current character and move cursor forward.
+ """
+ buffer = event.current_buffer
+ c = buffer.document.current_char
+
+ if c is not None and c != "\n":
+ buffer.insert_text(c.swapcase(), overwrite=True)
+
+ @handle("g", "u", "u", filter=vi_navigation_mode & ~is_read_only)
+ def _lowercase_line(event: E) -> None:
+ """
+ Lowercase current line.
+ """
+ buff = event.current_buffer
+ buff.transform_current_line(lambda s: s.lower())
+
+ @handle("g", "U", "U", filter=vi_navigation_mode & ~is_read_only)
+ def _uppercase_line(event: E) -> None:
+ """
+ Uppercase current line.
+ """
+ buff = event.current_buffer
+ buff.transform_current_line(lambda s: s.upper())
+
+ @handle("g", "~", "~", filter=vi_navigation_mode & ~is_read_only)
+ def _swapcase_line(event: E) -> None:
+ """
+ Swap case of the current line.
+ """
+ buff = event.current_buffer
+ buff.transform_current_line(lambda s: s.swapcase())
+
+ @handle("#", filter=vi_navigation_mode)
+ def _prev_occurrence(event: E) -> None:
+ """
+ Go to previous occurrence of this word.
+ """
+ b = event.current_buffer
+ search_state = event.app.current_search_state
+
+ search_state.text = b.document.get_word_under_cursor()
+ search_state.direction = SearchDirection.BACKWARD
+
+ b.apply_search(search_state, count=event.arg, include_current_position=False)
+
+ @handle("*", filter=vi_navigation_mode)
+ def _next_occurrence(event: E) -> None:
+ """
+ Go to next occurrence of this word.
+ """
+ b = event.current_buffer
+ search_state = event.app.current_search_state
+
+ search_state.text = b.document.get_word_under_cursor()
+ search_state.direction = SearchDirection.FORWARD
+
+ b.apply_search(search_state, count=event.arg, include_current_position=False)
+
+ @handle("(", filter=vi_navigation_mode)
+ def _begin_of_sentence(event: E) -> None:
+ # TODO: go to begin of sentence.
+ # XXX: should become text_object.
+ pass
+
+ @handle(")", filter=vi_navigation_mode)
+ def _end_of_sentence(event: E) -> None:
+ # TODO: go to end of sentence.
+ # XXX: should become text_object.
+ pass
+
+ operator = create_operator_decorator(key_bindings)
+ text_object = create_text_object_decorator(key_bindings)
+
+ @handle(Keys.Any, filter=vi_waiting_for_text_object_mode)
+ def _unknown_text_object(event: E) -> None:
+ """
+ Unknown key binding while waiting for a text object.
+ """
+ event.app.output.bell()
+
+ #
+ # *** Operators ***
+ #
+
+ def create_delete_and_change_operators(
+ delete_only: bool, with_register: bool = False
+ ) -> None:
+ """
+ Delete and change operators.
+
+ :param delete_only: Create an operator that deletes, but doesn't go to insert mode.
+ :param with_register: Copy the deleted text to this named register instead of the clipboard.
+ """
+ handler_keys: Iterable[str]
+ if with_register:
+ handler_keys = ('"', Keys.Any, "cd"[delete_only])
+ else:
+ handler_keys = "cd"[delete_only]
+
+ @operator(*handler_keys, filter=~is_read_only)
+ def delete_or_change_operator(event: E, text_object: TextObject) -> None:
+ clipboard_data = None
+ buff = event.current_buffer
+
+ if text_object:
+ new_document, clipboard_data = text_object.cut(buff)
+ buff.document = new_document
+
+ # Set deleted/changed text to clipboard or named register.
+ if clipboard_data and clipboard_data.text:
+ if with_register:
+ reg_name = event.key_sequence[1].data
+ if reg_name in vi_register_names:
+ event.app.vi_state.named_registers[reg_name] = clipboard_data
+ else:
+ event.app.clipboard.set_data(clipboard_data)
+
+ # Only go back to insert mode in case of 'change'.
+ if not delete_only:
+ event.app.vi_state.input_mode = InputMode.INSERT
+
+ create_delete_and_change_operators(False, False)
+ create_delete_and_change_operators(False, True)
+ create_delete_and_change_operators(True, False)
+ create_delete_and_change_operators(True, True)
+
+ def create_transform_handler(
+ filter: Filter, transform_func: Callable[[str], str], *a: str
+ ) -> None:
+ @operator(*a, filter=filter & ~is_read_only)
+ def _(event: E, text_object: TextObject) -> None:
+ """
+ Apply transformation (uppercase, lowercase, rot13, swap case).
+ """
+ buff = event.current_buffer
+ start, end = text_object.operator_range(buff.document)
+
+ if start < end:
+ # Transform.
+ buff.transform_region(
+ buff.cursor_position + start,
+ buff.cursor_position + end,
+ transform_func,
+ )
+
+ # Move cursor
+ buff.cursor_position += text_object.end or text_object.start
+
+ for k, f, func in vi_transform_functions:
+ create_transform_handler(f, func, *k)
+
+ @operator("y")
+ def _yank(event: E, text_object: TextObject) -> None:
+ """
+ Yank operator. (Copy text.)
+ """
+ _, clipboard_data = text_object.cut(event.current_buffer)
+ if clipboard_data.text:
+ event.app.clipboard.set_data(clipboard_data)
+
+ @operator('"', Keys.Any, "y")
+ def _yank_to_register(event: E, text_object: TextObject) -> None:
+ """
+ Yank selection to named register.
+ """
+ c = event.key_sequence[1].data
+ if c in vi_register_names:
+ _, clipboard_data = text_object.cut(event.current_buffer)
+ event.app.vi_state.named_registers[c] = clipboard_data
+
+ @operator(">")
+ def _indent_text_object(event: E, text_object: TextObject) -> None:
+ """
+ Indent.
+ """
+ buff = event.current_buffer
+ from_, to = text_object.get_line_numbers(buff)
+ indent(buff, from_, to + 1, count=event.arg)
+
+ @operator("<")
+ def _unindent_text_object(event: E, text_object: TextObject) -> None:
+ """
+ Unindent.
+ """
+ buff = event.current_buffer
+ from_, to = text_object.get_line_numbers(buff)
+ unindent(buff, from_, to + 1, count=event.arg)
+
+ @operator("g", "q")
+ def _reshape(event: E, text_object: TextObject) -> None:
+ """
+ Reshape text.
+ """
+ buff = event.current_buffer
+ from_, to = text_object.get_line_numbers(buff)
+ reshape_text(buff, from_, to)
+
+ #
+ # *** Text objects ***
+ #
+
+ @text_object("b")
+ def _b(event: E) -> TextObject:
+ """
+ Move one word or token left.
+ """
+ return TextObject(
+ event.current_buffer.document.find_start_of_previous_word(count=event.arg)
+ or 0
+ )
+
+ @text_object("B")
+ def _B(event: E) -> TextObject:
+ """
+ Move one non-blank word left
+ """
+ return TextObject(
+ event.current_buffer.document.find_start_of_previous_word(
+ count=event.arg, WORD=True
+ )
+ or 0
+ )
+
+ @text_object("$")
+ def _dollar(event: E) -> TextObject:
+ """
+ 'c$', 'd$' and '$': Delete/change/move until end of line.
+ """
+ return TextObject(event.current_buffer.document.get_end_of_line_position())
+
+ @text_object("w")
+ def _word_forward(event: E) -> TextObject:
+ """
+ 'word' forward. 'cw', 'dw', 'w': Delete/change/move one word.
+ """
+ return TextObject(
+ event.current_buffer.document.find_next_word_beginning(count=event.arg)
+ or event.current_buffer.document.get_end_of_document_position()
+ )
+
+ @text_object("W")
+ def _WORD_forward(event: E) -> TextObject:
+ """
+ 'WORD' forward. 'cW', 'dW', 'W': Delete/change/move one WORD.
+ """
+ return TextObject(
+ event.current_buffer.document.find_next_word_beginning(
+ count=event.arg, WORD=True
+ )
+ or event.current_buffer.document.get_end_of_document_position()
+ )
+
+ @text_object("e")
+ def _end_of_word(event: E) -> TextObject:
+ """
+ End of 'word': 'ce', 'de', 'e'
+ """
+ end = event.current_buffer.document.find_next_word_ending(count=event.arg)
+ return TextObject(end - 1 if end else 0, type=TextObjectType.INCLUSIVE)
+
+ @text_object("E")
+ def _end_of_WORD(event: E) -> TextObject:
+ """
+ End of 'WORD': 'cE', 'dE', 'E'
+ """
+ end = event.current_buffer.document.find_next_word_ending(
+ count=event.arg, WORD=True
+ )
+ return TextObject(end - 1 if end else 0, type=TextObjectType.INCLUSIVE)
+
+ @text_object("i", "w", no_move_handler=True)
+ def _inner_word(event: E) -> TextObject:
+ """
+ Inner 'word': ciw and diw
+ """
+ start, end = event.current_buffer.document.find_boundaries_of_current_word()
+ return TextObject(start, end)
+
+ @text_object("a", "w", no_move_handler=True)
+ def _a_word(event: E) -> TextObject:
+ """
+ A 'word': caw and daw
+ """
+ start, end = event.current_buffer.document.find_boundaries_of_current_word(
+ include_trailing_whitespace=True
+ )
+ return TextObject(start, end)
+
+ @text_object("i", "W", no_move_handler=True)
+ def _inner_WORD(event: E) -> TextObject:
+ """
+ Inner 'WORD': ciW and diW
+ """
+ start, end = event.current_buffer.document.find_boundaries_of_current_word(
+ WORD=True
+ )
+ return TextObject(start, end)
+
+ @text_object("a", "W", no_move_handler=True)
+ def _a_WORD(event: E) -> TextObject:
+ """
+ A 'WORD': caw and daw
+ """
+ start, end = event.current_buffer.document.find_boundaries_of_current_word(
+ WORD=True, include_trailing_whitespace=True
+ )
+ return TextObject(start, end)
+
+ @text_object("a", "p", no_move_handler=True)
+ def _paragraph(event: E) -> TextObject:
+ """
+ Auto paragraph.
+ """
+ start = event.current_buffer.document.start_of_paragraph()
+ end = event.current_buffer.document.end_of_paragraph(count=event.arg)
+ return TextObject(start, end)
+
+ @text_object("^")
+ def _start_of_line(event: E) -> TextObject:
+ """'c^', 'd^' and '^': Soft start of line, after whitespace."""
+ return TextObject(
+ event.current_buffer.document.get_start_of_line_position(
+ after_whitespace=True
+ )
+ )
+
+ @text_object("0")
+ def _hard_start_of_line(event: E) -> TextObject:
+ """
+ 'c0', 'd0': Hard start of line, before whitespace.
+ (The move '0' key is implemented elsewhere, because a '0' could also change the `arg`.)
+ """
+ return TextObject(
+ event.current_buffer.document.get_start_of_line_position(
+ after_whitespace=False
+ )
+ )
+
+ def create_ci_ca_handles(
+ ci_start: str, ci_end: str, inner: bool, key: str | None = None
+ ) -> None:
+ # TODO: 'dat', 'dit', (tags (like xml)
+ """
+ Delete/Change string between this start and stop character. But keep these characters.
+ This implements all the ci", ci<, ci{, ci(, di", di<, ca", ca<, ... combinations.
+ """
+
+ def handler(event: E) -> TextObject:
+ if ci_start == ci_end:
+ # Quotes
+ start = event.current_buffer.document.find_backwards(
+ ci_start, in_current_line=False
+ )
+ end = event.current_buffer.document.find(ci_end, in_current_line=False)
+ else:
+ # Brackets
+ start = event.current_buffer.document.find_enclosing_bracket_left(
+ ci_start, ci_end
+ )
+ end = event.current_buffer.document.find_enclosing_bracket_right(
+ ci_start, ci_end
+ )
+
+ if start is not None and end is not None:
+ offset = 0 if inner else 1
+ return TextObject(start + 1 - offset, end + offset)
+ else:
+ # Nothing found.
+ return TextObject(0)
+
+ if key is None:
+ text_object("ai"[inner], ci_start, no_move_handler=True)(handler)
+ text_object("ai"[inner], ci_end, no_move_handler=True)(handler)
+ else:
+ text_object("ai"[inner], key, no_move_handler=True)(handler)
+
+ for inner in (False, True):
+ for ci_start, ci_end in [
+ ('"', '"'),
+ ("'", "'"),
+ ("`", "`"),
+ ("[", "]"),
+ ("<", ">"),
+ ("{", "}"),
+ ("(", ")"),
+ ]:
+ create_ci_ca_handles(ci_start, ci_end, inner)
+
+ create_ci_ca_handles("(", ")", inner, "b") # 'dab', 'dib'
+ create_ci_ca_handles("{", "}", inner, "B") # 'daB', 'diB'
+
+ @text_object("{")
+ def _previous_section(event: E) -> TextObject:
+ """
+ Move to previous blank-line separated section.
+ Implements '{', 'c{', 'd{', 'y{'
+ """
+ index = event.current_buffer.document.start_of_paragraph(
+ count=event.arg, before=True
+ )
+ return TextObject(index)
+
+ @text_object("}")
+ def _next_section(event: E) -> TextObject:
+ """
+ Move to next blank-line separated section.
+ Implements '}', 'c}', 'd}', 'y}'
+ """
+ index = event.current_buffer.document.end_of_paragraph(
+ count=event.arg, after=True
+ )
+ return TextObject(index)
+
+ @text_object("f", Keys.Any)
+ def _find_next_occurrence(event: E) -> TextObject:
+ """
+ Go to next occurrence of character. Typing 'fx' will move the
+ cursor to the next occurrence of character. 'x'.
+ """
+ event.app.vi_state.last_character_find = CharacterFind(event.data, False)
+ match = event.current_buffer.document.find(
+ event.data, in_current_line=True, count=event.arg
+ )
+ if match:
+ return TextObject(match, type=TextObjectType.INCLUSIVE)
+ else:
+ return TextObject(0)
+
+ @text_object("F", Keys.Any)
+ def _find_previous_occurrence(event: E) -> TextObject:
+ """
+ Go to previous occurrence of character. Typing 'Fx' will move the
+ cursor to the previous occurrence of character. 'x'.
+ """
+ event.app.vi_state.last_character_find = CharacterFind(event.data, True)
+ return TextObject(
+ event.current_buffer.document.find_backwards(
+ event.data, in_current_line=True, count=event.arg
+ )
+ or 0
+ )
+
+ @text_object("t", Keys.Any)
+ def _t(event: E) -> TextObject:
+ """
+ Move right to the next occurrence of c, then one char backward.
+ """
+ event.app.vi_state.last_character_find = CharacterFind(event.data, False)
+ match = event.current_buffer.document.find(
+ event.data, in_current_line=True, count=event.arg
+ )
+ if match:
+ return TextObject(match - 1, type=TextObjectType.INCLUSIVE)
+ else:
+ return TextObject(0)
+
+ @text_object("T", Keys.Any)
+ def _T(event: E) -> TextObject:
+ """
+ Move left to the previous occurrence of c, then one char forward.
+ """
+ event.app.vi_state.last_character_find = CharacterFind(event.data, True)
+ match = event.current_buffer.document.find_backwards(
+ event.data, in_current_line=True, count=event.arg
+ )
+ return TextObject(match + 1 if match else 0)
+
+ def repeat(reverse: bool) -> None:
+ """
+ Create ',' and ';' commands.
+ """
+
+ @text_object("," if reverse else ";")
+ def _(event: E) -> TextObject:
+ """
+ Repeat the last 'f'/'F'/'t'/'T' command.
+ """
+ pos: int | None = 0
+ vi_state = event.app.vi_state
+
+ type = TextObjectType.EXCLUSIVE
+
+ if vi_state.last_character_find:
+ char = vi_state.last_character_find.character
+ backwards = vi_state.last_character_find.backwards
+
+ if reverse:
+ backwards = not backwards
+
+ if backwards:
+ pos = event.current_buffer.document.find_backwards(
+ char, in_current_line=True, count=event.arg
+ )
+ else:
+ pos = event.current_buffer.document.find(
+ char, in_current_line=True, count=event.arg
+ )
+ type = TextObjectType.INCLUSIVE
+ if pos:
+ return TextObject(pos, type=type)
+ else:
+ return TextObject(0)
+
+ repeat(True)
+ repeat(False)
+
+ @text_object("h")
+ @text_object("left")
+ def _left(event: E) -> TextObject:
+ """
+ Implements 'ch', 'dh', 'h': Cursor left.
+ """
+ return TextObject(
+ event.current_buffer.document.get_cursor_left_position(count=event.arg)
+ )
+
+ @text_object("j", no_move_handler=True, no_selection_handler=True)
+ # Note: We also need `no_selection_handler`, because we in
+ # selection mode, we prefer the other 'j' binding that keeps
+ # `buffer.preferred_column`.
+ def _down(event: E) -> TextObject:
+ """
+ Implements 'cj', 'dj', 'j', ... Cursor up.
+ """
+ return TextObject(
+ event.current_buffer.document.get_cursor_down_position(count=event.arg),
+ type=TextObjectType.LINEWISE,
+ )
+
+ @text_object("k", no_move_handler=True, no_selection_handler=True)
+ def _up(event: E) -> TextObject:
+ """
+ Implements 'ck', 'dk', 'k', ... Cursor up.
+ """
+ return TextObject(
+ event.current_buffer.document.get_cursor_up_position(count=event.arg),
+ type=TextObjectType.LINEWISE,
+ )
+
+ @text_object("l")
+ @text_object(" ")
+ @text_object("right")
+ def _right(event: E) -> TextObject:
+ """
+ Implements 'cl', 'dl', 'l', 'c ', 'd ', ' '. Cursor right.
+ """
+ return TextObject(
+ event.current_buffer.document.get_cursor_right_position(count=event.arg)
+ )
+
+ @text_object("H")
+ def _top_of_screen(event: E) -> TextObject:
+ """
+ Moves to the start of the visible region. (Below the scroll offset.)
+ Implements 'cH', 'dH', 'H'.
+ """
+ w = event.app.layout.current_window
+ b = event.current_buffer
+
+ if w and w.render_info:
+ # When we find a Window that has BufferControl showing this window,
+ # move to the start of the visible area.
+ pos = (
+ b.document.translate_row_col_to_index(
+ w.render_info.first_visible_line(after_scroll_offset=True), 0
+ )
+ - b.cursor_position
+ )
+
+ else:
+ # Otherwise, move to the start of the input.
+ pos = -len(b.document.text_before_cursor)
+ return TextObject(pos, type=TextObjectType.LINEWISE)
+
+ @text_object("M")
+ def _middle_of_screen(event: E) -> TextObject:
+ """
+ Moves cursor to the vertical center of the visible region.
+ Implements 'cM', 'dM', 'M'.
+ """
+ w = event.app.layout.current_window
+ b = event.current_buffer
+
+ if w and w.render_info:
+ # When we find a Window that has BufferControl showing this window,
+ # move to the center of the visible area.
+ pos = (
+ b.document.translate_row_col_to_index(
+ w.render_info.center_visible_line(), 0
+ )
+ - b.cursor_position
+ )
+
+ else:
+ # Otherwise, move to the start of the input.
+ pos = -len(b.document.text_before_cursor)
+ return TextObject(pos, type=TextObjectType.LINEWISE)
+
+ @text_object("L")
+ def _end_of_screen(event: E) -> TextObject:
+ """
+ Moves to the end of the visible region. (Above the scroll offset.)
+ """
+ w = event.app.layout.current_window
+ b = event.current_buffer
+
+ if w and w.render_info:
+ # When we find a Window that has BufferControl showing this window,
+ # move to the end of the visible area.
+ pos = (
+ b.document.translate_row_col_to_index(
+ w.render_info.last_visible_line(before_scroll_offset=True), 0
+ )
+ - b.cursor_position
+ )
+
+ else:
+ # Otherwise, move to the end of the input.
+ pos = len(b.document.text_after_cursor)
+ return TextObject(pos, type=TextObjectType.LINEWISE)
+
+ @text_object("n", no_move_handler=True)
+ def _search_next(event: E) -> TextObject:
+ """
+ Search next.
+ """
+ buff = event.current_buffer
+ search_state = event.app.current_search_state
+
+ cursor_position = buff.get_search_position(
+ search_state, include_current_position=False, count=event.arg
+ )
+ return TextObject(cursor_position - buff.cursor_position)
+
+ @handle("n", filter=vi_navigation_mode)
+ def _search_next2(event: E) -> None:
+ """
+ Search next in navigation mode. (This goes through the history.)
+ """
+ search_state = event.app.current_search_state
+
+ event.current_buffer.apply_search(
+ search_state, include_current_position=False, count=event.arg
+ )
+
+ @text_object("N", no_move_handler=True)
+ def _search_previous(event: E) -> TextObject:
+ """
+ Search previous.
+ """
+ buff = event.current_buffer
+ search_state = event.app.current_search_state
+
+ cursor_position = buff.get_search_position(
+ ~search_state, include_current_position=False, count=event.arg
+ )
+ return TextObject(cursor_position - buff.cursor_position)
+
+ @handle("N", filter=vi_navigation_mode)
+ def _search_previous2(event: E) -> None:
+ """
+ Search previous in navigation mode. (This goes through the history.)
+ """
+ search_state = event.app.current_search_state
+
+ event.current_buffer.apply_search(
+ ~search_state, include_current_position=False, count=event.arg
+ )
+
+ @handle("z", "+", filter=vi_navigation_mode | vi_selection_mode)
+ @handle("z", "t", filter=vi_navigation_mode | vi_selection_mode)
+ @handle("z", "enter", filter=vi_navigation_mode | vi_selection_mode)
+ def _scroll_top(event: E) -> None:
+ """
+ Scrolls the window to makes the current line the first line in the visible region.
+ """
+ b = event.current_buffer
+ event.app.layout.current_window.vertical_scroll = b.document.cursor_position_row
+
+ @handle("z", "-", filter=vi_navigation_mode | vi_selection_mode)
+ @handle("z", "b", filter=vi_navigation_mode | vi_selection_mode)
+ def _scroll_bottom(event: E) -> None:
+ """
+ Scrolls the window to makes the current line the last line in the visible region.
+ """
+ # We can safely set the scroll offset to zero; the Window will make
+ # sure that it scrolls at least enough to make the cursor visible
+ # again.
+ event.app.layout.current_window.vertical_scroll = 0
+
+ @handle("z", "z", filter=vi_navigation_mode | vi_selection_mode)
+ def _scroll_center(event: E) -> None:
+ """
+ Center Window vertically around cursor.
+ """
+ w = event.app.layout.current_window
+ b = event.current_buffer
+
+ if w and w.render_info:
+ info = w.render_info
+
+ # Calculate the offset that we need in order to position the row
+ # containing the cursor in the center.
+ scroll_height = info.window_height // 2
+
+ y = max(0, b.document.cursor_position_row - 1)
+ height = 0
+ while y > 0:
+ line_height = info.get_height_for_line(y)
+
+ if height + line_height < scroll_height:
+ height += line_height
+ y -= 1
+ else:
+ break
+
+ w.vertical_scroll = y
+
+ @text_object("%")
+ def _goto_corresponding_bracket(event: E) -> TextObject:
+ """
+ Implements 'c%', 'd%', '%, 'y%' (Move to corresponding bracket.)
+ If an 'arg' has been given, go this this % position in the file.
+ """
+ buffer = event.current_buffer
+
+ if event._arg:
+ # If 'arg' has been given, the meaning of % is to go to the 'x%'
+ # row in the file.
+ if 0 < event.arg <= 100:
+ absolute_index = buffer.document.translate_row_col_to_index(
+ int((event.arg * buffer.document.line_count - 1) / 100), 0
+ )
+ return TextObject(
+ absolute_index - buffer.document.cursor_position,
+ type=TextObjectType.LINEWISE,
+ )
+ else:
+ return TextObject(0) # Do nothing.
+
+ else:
+ # Move to the corresponding opening/closing bracket (()'s, []'s and {}'s).
+ match = buffer.document.find_matching_bracket_position()
+ if match:
+ return TextObject(match, type=TextObjectType.INCLUSIVE)
+ else:
+ return TextObject(0)
+
+ @text_object("|")
+ def _to_column(event: E) -> TextObject:
+ """
+ Move to the n-th column (you may specify the argument n by typing it on
+ number keys, for example, 20|).
+ """
+ return TextObject(
+ event.current_buffer.document.get_column_cursor_position(event.arg - 1)
+ )
+
+ @text_object("g", "g")
+ def _goto_first_line(event: E) -> TextObject:
+ """
+ Go to the start of the very first line.
+ Implements 'gg', 'cgg', 'ygg'
+ """
+ d = event.current_buffer.document
+
+ if event._arg:
+ # Move to the given line.
+ return TextObject(
+ d.translate_row_col_to_index(event.arg - 1, 0) - d.cursor_position,
+ type=TextObjectType.LINEWISE,
+ )
+ else:
+ # Move to the top of the input.
+ return TextObject(
+ d.get_start_of_document_position(), type=TextObjectType.LINEWISE
+ )
+
+ @text_object("g", "_")
+ def _goto_last_line(event: E) -> TextObject:
+ """
+ Go to last non-blank of line.
+ 'g_', 'cg_', 'yg_', etc..
+ """
+ return TextObject(
+ event.current_buffer.document.last_non_blank_of_current_line_position(),
+ type=TextObjectType.INCLUSIVE,
+ )
+
+ @text_object("g", "e")
+ def _ge(event: E) -> TextObject:
+ """
+ Go to last character of previous word.
+ 'ge', 'cge', 'yge', etc..
+ """
+ prev_end = event.current_buffer.document.find_previous_word_ending(
+ count=event.arg
+ )
+ return TextObject(
+ prev_end - 1 if prev_end is not None else 0, type=TextObjectType.INCLUSIVE
+ )
+
+ @text_object("g", "E")
+ def _gE(event: E) -> TextObject:
+ """
+ Go to last character of previous WORD.
+ 'gE', 'cgE', 'ygE', etc..
+ """
+ prev_end = event.current_buffer.document.find_previous_word_ending(
+ count=event.arg, WORD=True
+ )
+ return TextObject(
+ prev_end - 1 if prev_end is not None else 0, type=TextObjectType.INCLUSIVE
+ )
+
+ @text_object("g", "m")
+ def _gm(event: E) -> TextObject:
+ """
+ Like g0, but half a screenwidth to the right. (Or as much as possible.)
+ """
+ w = event.app.layout.current_window
+ buff = event.current_buffer
+
+ if w and w.render_info:
+ width = w.render_info.window_width
+ start = buff.document.get_start_of_line_position(after_whitespace=False)
+ start += int(min(width / 2, len(buff.document.current_line)))
+
+ return TextObject(start, type=TextObjectType.INCLUSIVE)
+ return TextObject(0)
+
+ @text_object("G")
+ def _last_line(event: E) -> TextObject:
+ """
+ Go to the end of the document. (If no arg has been given.)
+ """
+ buf = event.current_buffer
+ return TextObject(
+ buf.document.translate_row_col_to_index(buf.document.line_count - 1, 0)
+ - buf.cursor_position,
+ type=TextObjectType.LINEWISE,
+ )
+
+ #
+ # *** Other ***
+ #
+
+ @handle("G", filter=has_arg)
+ def _to_nth_history_line(event: E) -> None:
+ """
+ If an argument is given, move to this line in the history. (for
+ example, 15G)
+ """
+ event.current_buffer.go_to_history(event.arg - 1)
+
+ for n in "123456789":
+
+ @handle(
+ n,
+ filter=vi_navigation_mode
+ | vi_selection_mode
+ | vi_waiting_for_text_object_mode,
+ )
+ def _arg(event: E) -> None:
+ """
+ Always handle numerics in navigation mode as arg.
+ """
+ event.append_to_arg_count(event.data)
+
+ @handle(
+ "0",
+ filter=(
+ vi_navigation_mode | vi_selection_mode | vi_waiting_for_text_object_mode
+ )
+ & has_arg,
+ )
+ def _0_arg(event: E) -> None:
+ """
+ Zero when an argument was already give.
+ """
+ event.append_to_arg_count(event.data)
+
+ @handle(Keys.Any, filter=vi_replace_mode)
+ def _insert_text(event: E) -> None:
+ """
+ Insert data at cursor position.
+ """
+ event.current_buffer.insert_text(event.data, overwrite=True)
+
+ @handle(Keys.Any, filter=vi_replace_single_mode)
+ def _replace_single(event: E) -> None:
+ """
+ Replace single character at cursor position.
+ """
+ event.current_buffer.insert_text(event.data, overwrite=True)
+ event.current_buffer.cursor_position -= 1
+ event.app.vi_state.input_mode = InputMode.NAVIGATION
+
+ @handle(
+ Keys.Any,
+ filter=vi_insert_multiple_mode,
+ save_before=(lambda e: not e.is_repeat),
+ )
+ def _insert_text_multiple_cursors(event: E) -> None:
+ """
+ Insert data at multiple cursor positions at once.
+ (Usually a result of pressing 'I' or 'A' in block-selection mode.)
+ """
+ buff = event.current_buffer
+ original_text = buff.text
+
+ # Construct new text.
+ text = []
+ p = 0
+
+ for p2 in buff.multiple_cursor_positions:
+ text.append(original_text[p:p2])
+ text.append(event.data)
+ p = p2
+
+ text.append(original_text[p:])
+
+ # Shift all cursor positions.
+ new_cursor_positions = [
+ pos + i + 1 for i, pos in enumerate(buff.multiple_cursor_positions)
+ ]
+
+ # Set result.
+ buff.text = "".join(text)
+ buff.multiple_cursor_positions = new_cursor_positions
+ buff.cursor_position += 1
+
+ @handle("backspace", filter=vi_insert_multiple_mode)
+ def _delete_before_multiple_cursors(event: E) -> None:
+ """
+ Backspace, using multiple cursors.
+ """
+ buff = event.current_buffer
+ original_text = buff.text
+
+ # Construct new text.
+ deleted_something = False
+ text = []
+ p = 0
+
+ for p2 in buff.multiple_cursor_positions:
+ if p2 > 0 and original_text[p2 - 1] != "\n": # Don't delete across lines.
+ text.append(original_text[p : p2 - 1])
+ deleted_something = True
+ else:
+ text.append(original_text[p:p2])
+ p = p2
+
+ text.append(original_text[p:])
+
+ if deleted_something:
+ # Shift all cursor positions.
+ lengths = [len(part) for part in text[:-1]]
+ new_cursor_positions = list(accumulate(lengths))
+
+ # Set result.
+ buff.text = "".join(text)
+ buff.multiple_cursor_positions = new_cursor_positions
+ buff.cursor_position -= 1
+ else:
+ event.app.output.bell()
+
+ @handle("delete", filter=vi_insert_multiple_mode)
+ def _delete_after_multiple_cursors(event: E) -> None:
+ """
+ Delete, using multiple cursors.
+ """
+ buff = event.current_buffer
+ original_text = buff.text
+
+ # Construct new text.
+ deleted_something = False
+ text = []
+ new_cursor_positions = []
+ p = 0
+
+ for p2 in buff.multiple_cursor_positions:
+ text.append(original_text[p:p2])
+ if p2 >= len(original_text) or original_text[p2] == "\n":
+ # Don't delete across lines.
+ p = p2
+ else:
+ p = p2 + 1
+ deleted_something = True
+
+ text.append(original_text[p:])
+
+ if deleted_something:
+ # Shift all cursor positions.
+ lengths = [len(part) for part in text[:-1]]
+ new_cursor_positions = list(accumulate(lengths))
+
+ # Set result.
+ buff.text = "".join(text)
+ buff.multiple_cursor_positions = new_cursor_positions
+ else:
+ event.app.output.bell()
+
+ @handle("left", filter=vi_insert_multiple_mode)
+ def _left_multiple(event: E) -> None:
+ """
+ Move all cursors to the left.
+ (But keep all cursors on the same line.)
+ """
+ buff = event.current_buffer
+ new_positions = []
+
+ for p in buff.multiple_cursor_positions:
+ if buff.document.translate_index_to_position(p)[1] > 0:
+ p -= 1
+ new_positions.append(p)
+
+ buff.multiple_cursor_positions = new_positions
+
+ if buff.document.cursor_position_col > 0:
+ buff.cursor_position -= 1
+
+ @handle("right", filter=vi_insert_multiple_mode)
+ def _right_multiple(event: E) -> None:
+ """
+ Move all cursors to the right.
+ (But keep all cursors on the same line.)
+ """
+ buff = event.current_buffer
+ new_positions = []
+
+ for p in buff.multiple_cursor_positions:
+ row, column = buff.document.translate_index_to_position(p)
+ if column < len(buff.document.lines[row]):
+ p += 1
+ new_positions.append(p)
+
+ buff.multiple_cursor_positions = new_positions
+
+ if not buff.document.is_cursor_at_the_end_of_line:
+ buff.cursor_position += 1
+
+ @handle("up", filter=vi_insert_multiple_mode)
+ @handle("down", filter=vi_insert_multiple_mode)
+ def _updown_multiple(event: E) -> None:
+ """
+ Ignore all up/down key presses when in multiple cursor mode.
+ """
+
+ @handle("c-x", "c-l", filter=vi_insert_mode)
+ def _complete_line(event: E) -> None:
+ """
+ Pressing the ControlX - ControlL sequence in Vi mode does line
+ completion based on the other lines in the document and the history.
+ """
+ event.current_buffer.start_history_lines_completion()
+
+ @handle("c-x", "c-f", filter=vi_insert_mode)
+ def _complete_filename(event: E) -> None:
+ """
+ Complete file names.
+ """
+ # TODO
+ pass
+
+ @handle("c-k", filter=vi_insert_mode | vi_replace_mode)
+ def _digraph(event: E) -> None:
+ """
+ Go into digraph mode.
+ """
+ event.app.vi_state.waiting_for_digraph = True
+
+ @handle(Keys.Any, filter=vi_digraph_mode & ~digraph_symbol_1_given)
+ def _digraph1(event: E) -> None:
+ """
+ First digraph symbol.
+ """
+ event.app.vi_state.digraph_symbol1 = event.data
+
+ @handle(Keys.Any, filter=vi_digraph_mode & digraph_symbol_1_given)
+ def _create_digraph(event: E) -> None:
+ """
+ Insert digraph.
+ """
+ try:
+ # Lookup.
+ code: tuple[str, str] = (
+ event.app.vi_state.digraph_symbol1 or "",
+ event.data,
+ )
+ if code not in DIGRAPHS:
+ code = code[::-1] # Try reversing.
+ symbol = DIGRAPHS[code]
+ except KeyError:
+ # Unknown digraph.
+ event.app.output.bell()
+ else:
+ # Insert digraph.
+ overwrite = event.app.vi_state.input_mode == InputMode.REPLACE
+ event.current_buffer.insert_text(chr(symbol), overwrite=overwrite)
+ event.app.vi_state.waiting_for_digraph = False
+ finally:
+ event.app.vi_state.waiting_for_digraph = False
+ event.app.vi_state.digraph_symbol1 = None
+
+ @handle("c-o", filter=vi_insert_mode | vi_replace_mode)
+ def _quick_normal_mode(event: E) -> None:
+ """
+ Go into normal mode for one single action.
+ """
+ event.app.vi_state.temporary_navigation_mode = True
+
+ @handle("q", Keys.Any, filter=vi_navigation_mode & ~vi_recording_macro)
+ def _start_macro(event: E) -> None:
+ """
+ Start recording macro.
+ """
+ c = event.key_sequence[1].data
+ if c in vi_register_names:
+ vi_state = event.app.vi_state
+
+ vi_state.recording_register = c
+ vi_state.current_recording = ""
+
+ @handle("q", filter=vi_navigation_mode & vi_recording_macro)
+ def _stop_macro(event: E) -> None:
+ """
+ Stop recording macro.
+ """
+ vi_state = event.app.vi_state
+
+ # Store and stop recording.
+ if vi_state.recording_register:
+ vi_state.named_registers[vi_state.recording_register] = ClipboardData(
+ vi_state.current_recording
+ )
+ vi_state.recording_register = None
+ vi_state.current_recording = ""
+
+ @handle("@", Keys.Any, filter=vi_navigation_mode, record_in_macro=False)
+ def _execute_macro(event: E) -> None:
+ """
+ Execute macro.
+
+ Notice that we pass `record_in_macro=False`. This ensures that the `@x`
+ keys don't appear in the recording itself. This function inserts the
+ body of the called macro back into the KeyProcessor, so these keys will
+ be added later on to the macro of their handlers have
+ `record_in_macro=True`.
+ """
+ # Retrieve macro.
+ c = event.key_sequence[1].data
+ try:
+ macro = event.app.vi_state.named_registers[c]
+ except KeyError:
+ return
+
+ # Expand macro (which is a string in the register), in individual keys.
+ # Use vt100 parser for this.
+ keys: list[KeyPress] = []
+
+ parser = Vt100Parser(keys.append)
+ parser.feed(macro.text)
+ parser.flush()
+
+ # Now feed keys back to the input processor.
+ for _ in range(event.arg):
+ event.app.key_processor.feed_multiple(keys, first=True)
+
+ return ConditionalKeyBindings(key_bindings, vi_mode)
+
+
+def load_vi_search_bindings() -> KeyBindingsBase:
+ key_bindings = KeyBindings()
+ handle = key_bindings.add
+ from . import search
+
+ # Vi-style forward search.
+ handle(
+ "/",
+ filter=(vi_navigation_mode | vi_selection_mode) & ~vi_search_direction_reversed,
+ )(search.start_forward_incremental_search)
+ handle(
+ "?",
+ filter=(vi_navigation_mode | vi_selection_mode) & vi_search_direction_reversed,
+ )(search.start_forward_incremental_search)
+ handle("c-s")(search.start_forward_incremental_search)
+
+ # Vi-style backward search.
+ handle(
+ "?",
+ filter=(vi_navigation_mode | vi_selection_mode) & ~vi_search_direction_reversed,
+ )(search.start_reverse_incremental_search)
+ handle(
+ "/",
+ filter=(vi_navigation_mode | vi_selection_mode) & vi_search_direction_reversed,
+ )(search.start_reverse_incremental_search)
+ handle("c-r")(search.start_reverse_incremental_search)
+
+ # Apply the search. (At the / or ? prompt.)
+ handle("enter", filter=is_searching)(search.accept_search)
+
+ handle("c-r", filter=is_searching)(search.reverse_incremental_search)
+ handle("c-s", filter=is_searching)(search.forward_incremental_search)
+
+ handle("c-c")(search.abort_search)
+ handle("c-g")(search.abort_search)
+ handle("backspace", filter=search_buffer_is_empty)(search.abort_search)
+
+ # Handle escape. This should accept the search, just like readline.
+ # `abort_search` would be a meaningful alternative.
+ handle("escape")(search.accept_search)
+
+ return ConditionalKeyBindings(key_bindings, vi_mode)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/defaults.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/defaults.py
new file mode 100644
index 0000000000000000000000000000000000000000..6c265711603dfa2d151fa3f34d548b5cf69c2820
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/defaults.py
@@ -0,0 +1,63 @@
+"""
+Default key bindings.::
+
+ key_bindings = load_key_bindings()
+ app = Application(key_bindings=key_bindings)
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit.filters import buffer_has_focus
+from prompt_toolkit.key_binding.bindings.basic import load_basic_bindings
+from prompt_toolkit.key_binding.bindings.cpr import load_cpr_bindings
+from prompt_toolkit.key_binding.bindings.emacs import (
+ load_emacs_bindings,
+ load_emacs_search_bindings,
+ load_emacs_shift_selection_bindings,
+)
+from prompt_toolkit.key_binding.bindings.mouse import load_mouse_bindings
+from prompt_toolkit.key_binding.bindings.vi import (
+ load_vi_bindings,
+ load_vi_search_bindings,
+)
+from prompt_toolkit.key_binding.key_bindings import (
+ ConditionalKeyBindings,
+ KeyBindingsBase,
+ merge_key_bindings,
+)
+
+__all__ = [
+ "load_key_bindings",
+]
+
+
+def load_key_bindings() -> KeyBindingsBase:
+ """
+ Create a KeyBindings object that contains the default key bindings.
+ """
+ all_bindings = merge_key_bindings(
+ [
+ # Load basic bindings.
+ load_basic_bindings(),
+ # Load emacs bindings.
+ load_emacs_bindings(),
+ load_emacs_search_bindings(),
+ load_emacs_shift_selection_bindings(),
+ # Load Vi bindings.
+ load_vi_bindings(),
+ load_vi_search_bindings(),
+ ]
+ )
+
+ return merge_key_bindings(
+ [
+ # Make sure that the above key bindings are only active if the
+ # currently focused control is a `BufferControl`. For other controls, we
+ # don't want these key bindings to intervene. (This would break "ptterm"
+ # for instance, which handles 'Keys.Any' in the user control itself.)
+ ConditionalKeyBindings(all_bindings, buffer_has_focus),
+ # Active, even when no buffer has been focused.
+ load_mouse_bindings(),
+ load_cpr_bindings(),
+ ]
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/digraphs.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/digraphs.py
new file mode 100644
index 0000000000000000000000000000000000000000..f0152dc4cb05525ac4d4589e67402d7ba55ba1d4
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/digraphs.py
@@ -0,0 +1,1378 @@
+"""
+Vi Digraphs.
+This is a list of special characters that can be inserted in Vi insert mode by
+pressing Control-K followed by to normal characters.
+
+Taken from Neovim and translated to Python:
+https://raw.githubusercontent.com/neovim/neovim/master/src/nvim/digraph.c
+"""
+
+from __future__ import annotations
+
+__all__ = [
+ "DIGRAPHS",
+]
+
+# digraphs for Unicode from RFC1345
+# (also work for ISO-8859-1 aka latin1)
+DIGRAPHS: dict[tuple[str, str], int] = {
+ ("N", "U"): 0x00,
+ ("S", "H"): 0x01,
+ ("S", "X"): 0x02,
+ ("E", "X"): 0x03,
+ ("E", "T"): 0x04,
+ ("E", "Q"): 0x05,
+ ("A", "K"): 0x06,
+ ("B", "L"): 0x07,
+ ("B", "S"): 0x08,
+ ("H", "T"): 0x09,
+ ("L", "F"): 0x0A,
+ ("V", "T"): 0x0B,
+ ("F", "F"): 0x0C,
+ ("C", "R"): 0x0D,
+ ("S", "O"): 0x0E,
+ ("S", "I"): 0x0F,
+ ("D", "L"): 0x10,
+ ("D", "1"): 0x11,
+ ("D", "2"): 0x12,
+ ("D", "3"): 0x13,
+ ("D", "4"): 0x14,
+ ("N", "K"): 0x15,
+ ("S", "Y"): 0x16,
+ ("E", "B"): 0x17,
+ ("C", "N"): 0x18,
+ ("E", "M"): 0x19,
+ ("S", "B"): 0x1A,
+ ("E", "C"): 0x1B,
+ ("F", "S"): 0x1C,
+ ("G", "S"): 0x1D,
+ ("R", "S"): 0x1E,
+ ("U", "S"): 0x1F,
+ ("S", "P"): 0x20,
+ ("N", "b"): 0x23,
+ ("D", "O"): 0x24,
+ ("A", "t"): 0x40,
+ ("<", "("): 0x5B,
+ ("/", "/"): 0x5C,
+ (")", ">"): 0x5D,
+ ("'", ">"): 0x5E,
+ ("'", "!"): 0x60,
+ ("(", "!"): 0x7B,
+ ("!", "!"): 0x7C,
+ ("!", ")"): 0x7D,
+ ("'", "?"): 0x7E,
+ ("D", "T"): 0x7F,
+ ("P", "A"): 0x80,
+ ("H", "O"): 0x81,
+ ("B", "H"): 0x82,
+ ("N", "H"): 0x83,
+ ("I", "N"): 0x84,
+ ("N", "L"): 0x85,
+ ("S", "A"): 0x86,
+ ("E", "S"): 0x87,
+ ("H", "S"): 0x88,
+ ("H", "J"): 0x89,
+ ("V", "S"): 0x8A,
+ ("P", "D"): 0x8B,
+ ("P", "U"): 0x8C,
+ ("R", "I"): 0x8D,
+ ("S", "2"): 0x8E,
+ ("S", "3"): 0x8F,
+ ("D", "C"): 0x90,
+ ("P", "1"): 0x91,
+ ("P", "2"): 0x92,
+ ("T", "S"): 0x93,
+ ("C", "C"): 0x94,
+ ("M", "W"): 0x95,
+ ("S", "G"): 0x96,
+ ("E", "G"): 0x97,
+ ("S", "S"): 0x98,
+ ("G", "C"): 0x99,
+ ("S", "C"): 0x9A,
+ ("C", "I"): 0x9B,
+ ("S", "T"): 0x9C,
+ ("O", "C"): 0x9D,
+ ("P", "M"): 0x9E,
+ ("A", "C"): 0x9F,
+ ("N", "S"): 0xA0,
+ ("!", "I"): 0xA1,
+ ("C", "t"): 0xA2,
+ ("P", "d"): 0xA3,
+ ("C", "u"): 0xA4,
+ ("Y", "e"): 0xA5,
+ ("B", "B"): 0xA6,
+ ("S", "E"): 0xA7,
+ ("'", ":"): 0xA8,
+ ("C", "o"): 0xA9,
+ ("-", "a"): 0xAA,
+ ("<", "<"): 0xAB,
+ ("N", "O"): 0xAC,
+ ("-", "-"): 0xAD,
+ ("R", "g"): 0xAE,
+ ("'", "m"): 0xAF,
+ ("D", "G"): 0xB0,
+ ("+", "-"): 0xB1,
+ ("2", "S"): 0xB2,
+ ("3", "S"): 0xB3,
+ ("'", "'"): 0xB4,
+ ("M", "y"): 0xB5,
+ ("P", "I"): 0xB6,
+ (".", "M"): 0xB7,
+ ("'", ","): 0xB8,
+ ("1", "S"): 0xB9,
+ ("-", "o"): 0xBA,
+ (">", ">"): 0xBB,
+ ("1", "4"): 0xBC,
+ ("1", "2"): 0xBD,
+ ("3", "4"): 0xBE,
+ ("?", "I"): 0xBF,
+ ("A", "!"): 0xC0,
+ ("A", "'"): 0xC1,
+ ("A", ">"): 0xC2,
+ ("A", "?"): 0xC3,
+ ("A", ":"): 0xC4,
+ ("A", "A"): 0xC5,
+ ("A", "E"): 0xC6,
+ ("C", ","): 0xC7,
+ ("E", "!"): 0xC8,
+ ("E", "'"): 0xC9,
+ ("E", ">"): 0xCA,
+ ("E", ":"): 0xCB,
+ ("I", "!"): 0xCC,
+ ("I", "'"): 0xCD,
+ ("I", ">"): 0xCE,
+ ("I", ":"): 0xCF,
+ ("D", "-"): 0xD0,
+ ("N", "?"): 0xD1,
+ ("O", "!"): 0xD2,
+ ("O", "'"): 0xD3,
+ ("O", ">"): 0xD4,
+ ("O", "?"): 0xD5,
+ ("O", ":"): 0xD6,
+ ("*", "X"): 0xD7,
+ ("O", "/"): 0xD8,
+ ("U", "!"): 0xD9,
+ ("U", "'"): 0xDA,
+ ("U", ">"): 0xDB,
+ ("U", ":"): 0xDC,
+ ("Y", "'"): 0xDD,
+ ("T", "H"): 0xDE,
+ ("s", "s"): 0xDF,
+ ("a", "!"): 0xE0,
+ ("a", "'"): 0xE1,
+ ("a", ">"): 0xE2,
+ ("a", "?"): 0xE3,
+ ("a", ":"): 0xE4,
+ ("a", "a"): 0xE5,
+ ("a", "e"): 0xE6,
+ ("c", ","): 0xE7,
+ ("e", "!"): 0xE8,
+ ("e", "'"): 0xE9,
+ ("e", ">"): 0xEA,
+ ("e", ":"): 0xEB,
+ ("i", "!"): 0xEC,
+ ("i", "'"): 0xED,
+ ("i", ">"): 0xEE,
+ ("i", ":"): 0xEF,
+ ("d", "-"): 0xF0,
+ ("n", "?"): 0xF1,
+ ("o", "!"): 0xF2,
+ ("o", "'"): 0xF3,
+ ("o", ">"): 0xF4,
+ ("o", "?"): 0xF5,
+ ("o", ":"): 0xF6,
+ ("-", ":"): 0xF7,
+ ("o", "/"): 0xF8,
+ ("u", "!"): 0xF9,
+ ("u", "'"): 0xFA,
+ ("u", ">"): 0xFB,
+ ("u", ":"): 0xFC,
+ ("y", "'"): 0xFD,
+ ("t", "h"): 0xFE,
+ ("y", ":"): 0xFF,
+ ("A", "-"): 0x0100,
+ ("a", "-"): 0x0101,
+ ("A", "("): 0x0102,
+ ("a", "("): 0x0103,
+ ("A", ";"): 0x0104,
+ ("a", ";"): 0x0105,
+ ("C", "'"): 0x0106,
+ ("c", "'"): 0x0107,
+ ("C", ">"): 0x0108,
+ ("c", ">"): 0x0109,
+ ("C", "."): 0x010A,
+ ("c", "."): 0x010B,
+ ("C", "<"): 0x010C,
+ ("c", "<"): 0x010D,
+ ("D", "<"): 0x010E,
+ ("d", "<"): 0x010F,
+ ("D", "/"): 0x0110,
+ ("d", "/"): 0x0111,
+ ("E", "-"): 0x0112,
+ ("e", "-"): 0x0113,
+ ("E", "("): 0x0114,
+ ("e", "("): 0x0115,
+ ("E", "."): 0x0116,
+ ("e", "."): 0x0117,
+ ("E", ";"): 0x0118,
+ ("e", ";"): 0x0119,
+ ("E", "<"): 0x011A,
+ ("e", "<"): 0x011B,
+ ("G", ">"): 0x011C,
+ ("g", ">"): 0x011D,
+ ("G", "("): 0x011E,
+ ("g", "("): 0x011F,
+ ("G", "."): 0x0120,
+ ("g", "."): 0x0121,
+ ("G", ","): 0x0122,
+ ("g", ","): 0x0123,
+ ("H", ">"): 0x0124,
+ ("h", ">"): 0x0125,
+ ("H", "/"): 0x0126,
+ ("h", "/"): 0x0127,
+ ("I", "?"): 0x0128,
+ ("i", "?"): 0x0129,
+ ("I", "-"): 0x012A,
+ ("i", "-"): 0x012B,
+ ("I", "("): 0x012C,
+ ("i", "("): 0x012D,
+ ("I", ";"): 0x012E,
+ ("i", ";"): 0x012F,
+ ("I", "."): 0x0130,
+ ("i", "."): 0x0131,
+ ("I", "J"): 0x0132,
+ ("i", "j"): 0x0133,
+ ("J", ">"): 0x0134,
+ ("j", ">"): 0x0135,
+ ("K", ","): 0x0136,
+ ("k", ","): 0x0137,
+ ("k", "k"): 0x0138,
+ ("L", "'"): 0x0139,
+ ("l", "'"): 0x013A,
+ ("L", ","): 0x013B,
+ ("l", ","): 0x013C,
+ ("L", "<"): 0x013D,
+ ("l", "<"): 0x013E,
+ ("L", "."): 0x013F,
+ ("l", "."): 0x0140,
+ ("L", "/"): 0x0141,
+ ("l", "/"): 0x0142,
+ ("N", "'"): 0x0143,
+ ("n", "'"): 0x0144,
+ ("N", ","): 0x0145,
+ ("n", ","): 0x0146,
+ ("N", "<"): 0x0147,
+ ("n", "<"): 0x0148,
+ ("'", "n"): 0x0149,
+ ("N", "G"): 0x014A,
+ ("n", "g"): 0x014B,
+ ("O", "-"): 0x014C,
+ ("o", "-"): 0x014D,
+ ("O", "("): 0x014E,
+ ("o", "("): 0x014F,
+ ("O", '"'): 0x0150,
+ ("o", '"'): 0x0151,
+ ("O", "E"): 0x0152,
+ ("o", "e"): 0x0153,
+ ("R", "'"): 0x0154,
+ ("r", "'"): 0x0155,
+ ("R", ","): 0x0156,
+ ("r", ","): 0x0157,
+ ("R", "<"): 0x0158,
+ ("r", "<"): 0x0159,
+ ("S", "'"): 0x015A,
+ ("s", "'"): 0x015B,
+ ("S", ">"): 0x015C,
+ ("s", ">"): 0x015D,
+ ("S", ","): 0x015E,
+ ("s", ","): 0x015F,
+ ("S", "<"): 0x0160,
+ ("s", "<"): 0x0161,
+ ("T", ","): 0x0162,
+ ("t", ","): 0x0163,
+ ("T", "<"): 0x0164,
+ ("t", "<"): 0x0165,
+ ("T", "/"): 0x0166,
+ ("t", "/"): 0x0167,
+ ("U", "?"): 0x0168,
+ ("u", "?"): 0x0169,
+ ("U", "-"): 0x016A,
+ ("u", "-"): 0x016B,
+ ("U", "("): 0x016C,
+ ("u", "("): 0x016D,
+ ("U", "0"): 0x016E,
+ ("u", "0"): 0x016F,
+ ("U", '"'): 0x0170,
+ ("u", '"'): 0x0171,
+ ("U", ";"): 0x0172,
+ ("u", ";"): 0x0173,
+ ("W", ">"): 0x0174,
+ ("w", ">"): 0x0175,
+ ("Y", ">"): 0x0176,
+ ("y", ">"): 0x0177,
+ ("Y", ":"): 0x0178,
+ ("Z", "'"): 0x0179,
+ ("z", "'"): 0x017A,
+ ("Z", "."): 0x017B,
+ ("z", "."): 0x017C,
+ ("Z", "<"): 0x017D,
+ ("z", "<"): 0x017E,
+ ("O", "9"): 0x01A0,
+ ("o", "9"): 0x01A1,
+ ("O", "I"): 0x01A2,
+ ("o", "i"): 0x01A3,
+ ("y", "r"): 0x01A6,
+ ("U", "9"): 0x01AF,
+ ("u", "9"): 0x01B0,
+ ("Z", "/"): 0x01B5,
+ ("z", "/"): 0x01B6,
+ ("E", "D"): 0x01B7,
+ ("A", "<"): 0x01CD,
+ ("a", "<"): 0x01CE,
+ ("I", "<"): 0x01CF,
+ ("i", "<"): 0x01D0,
+ ("O", "<"): 0x01D1,
+ ("o", "<"): 0x01D2,
+ ("U", "<"): 0x01D3,
+ ("u", "<"): 0x01D4,
+ ("A", "1"): 0x01DE,
+ ("a", "1"): 0x01DF,
+ ("A", "7"): 0x01E0,
+ ("a", "7"): 0x01E1,
+ ("A", "3"): 0x01E2,
+ ("a", "3"): 0x01E3,
+ ("G", "/"): 0x01E4,
+ ("g", "/"): 0x01E5,
+ ("G", "<"): 0x01E6,
+ ("g", "<"): 0x01E7,
+ ("K", "<"): 0x01E8,
+ ("k", "<"): 0x01E9,
+ ("O", ";"): 0x01EA,
+ ("o", ";"): 0x01EB,
+ ("O", "1"): 0x01EC,
+ ("o", "1"): 0x01ED,
+ ("E", "Z"): 0x01EE,
+ ("e", "z"): 0x01EF,
+ ("j", "<"): 0x01F0,
+ ("G", "'"): 0x01F4,
+ ("g", "'"): 0x01F5,
+ (";", "S"): 0x02BF,
+ ("'", "<"): 0x02C7,
+ ("'", "("): 0x02D8,
+ ("'", "."): 0x02D9,
+ ("'", "0"): 0x02DA,
+ ("'", ";"): 0x02DB,
+ ("'", '"'): 0x02DD,
+ ("A", "%"): 0x0386,
+ ("E", "%"): 0x0388,
+ ("Y", "%"): 0x0389,
+ ("I", "%"): 0x038A,
+ ("O", "%"): 0x038C,
+ ("U", "%"): 0x038E,
+ ("W", "%"): 0x038F,
+ ("i", "3"): 0x0390,
+ ("A", "*"): 0x0391,
+ ("B", "*"): 0x0392,
+ ("G", "*"): 0x0393,
+ ("D", "*"): 0x0394,
+ ("E", "*"): 0x0395,
+ ("Z", "*"): 0x0396,
+ ("Y", "*"): 0x0397,
+ ("H", "*"): 0x0398,
+ ("I", "*"): 0x0399,
+ ("K", "*"): 0x039A,
+ ("L", "*"): 0x039B,
+ ("M", "*"): 0x039C,
+ ("N", "*"): 0x039D,
+ ("C", "*"): 0x039E,
+ ("O", "*"): 0x039F,
+ ("P", "*"): 0x03A0,
+ ("R", "*"): 0x03A1,
+ ("S", "*"): 0x03A3,
+ ("T", "*"): 0x03A4,
+ ("U", "*"): 0x03A5,
+ ("F", "*"): 0x03A6,
+ ("X", "*"): 0x03A7,
+ ("Q", "*"): 0x03A8,
+ ("W", "*"): 0x03A9,
+ ("J", "*"): 0x03AA,
+ ("V", "*"): 0x03AB,
+ ("a", "%"): 0x03AC,
+ ("e", "%"): 0x03AD,
+ ("y", "%"): 0x03AE,
+ ("i", "%"): 0x03AF,
+ ("u", "3"): 0x03B0,
+ ("a", "*"): 0x03B1,
+ ("b", "*"): 0x03B2,
+ ("g", "*"): 0x03B3,
+ ("d", "*"): 0x03B4,
+ ("e", "*"): 0x03B5,
+ ("z", "*"): 0x03B6,
+ ("y", "*"): 0x03B7,
+ ("h", "*"): 0x03B8,
+ ("i", "*"): 0x03B9,
+ ("k", "*"): 0x03BA,
+ ("l", "*"): 0x03BB,
+ ("m", "*"): 0x03BC,
+ ("n", "*"): 0x03BD,
+ ("c", "*"): 0x03BE,
+ ("o", "*"): 0x03BF,
+ ("p", "*"): 0x03C0,
+ ("r", "*"): 0x03C1,
+ ("*", "s"): 0x03C2,
+ ("s", "*"): 0x03C3,
+ ("t", "*"): 0x03C4,
+ ("u", "*"): 0x03C5,
+ ("f", "*"): 0x03C6,
+ ("x", "*"): 0x03C7,
+ ("q", "*"): 0x03C8,
+ ("w", "*"): 0x03C9,
+ ("j", "*"): 0x03CA,
+ ("v", "*"): 0x03CB,
+ ("o", "%"): 0x03CC,
+ ("u", "%"): 0x03CD,
+ ("w", "%"): 0x03CE,
+ ("'", "G"): 0x03D8,
+ (",", "G"): 0x03D9,
+ ("T", "3"): 0x03DA,
+ ("t", "3"): 0x03DB,
+ ("M", "3"): 0x03DC,
+ ("m", "3"): 0x03DD,
+ ("K", "3"): 0x03DE,
+ ("k", "3"): 0x03DF,
+ ("P", "3"): 0x03E0,
+ ("p", "3"): 0x03E1,
+ ("'", "%"): 0x03F4,
+ ("j", "3"): 0x03F5,
+ ("I", "O"): 0x0401,
+ ("D", "%"): 0x0402,
+ ("G", "%"): 0x0403,
+ ("I", "E"): 0x0404,
+ ("D", "S"): 0x0405,
+ ("I", "I"): 0x0406,
+ ("Y", "I"): 0x0407,
+ ("J", "%"): 0x0408,
+ ("L", "J"): 0x0409,
+ ("N", "J"): 0x040A,
+ ("T", "s"): 0x040B,
+ ("K", "J"): 0x040C,
+ ("V", "%"): 0x040E,
+ ("D", "Z"): 0x040F,
+ ("A", "="): 0x0410,
+ ("B", "="): 0x0411,
+ ("V", "="): 0x0412,
+ ("G", "="): 0x0413,
+ ("D", "="): 0x0414,
+ ("E", "="): 0x0415,
+ ("Z", "%"): 0x0416,
+ ("Z", "="): 0x0417,
+ ("I", "="): 0x0418,
+ ("J", "="): 0x0419,
+ ("K", "="): 0x041A,
+ ("L", "="): 0x041B,
+ ("M", "="): 0x041C,
+ ("N", "="): 0x041D,
+ ("O", "="): 0x041E,
+ ("P", "="): 0x041F,
+ ("R", "="): 0x0420,
+ ("S", "="): 0x0421,
+ ("T", "="): 0x0422,
+ ("U", "="): 0x0423,
+ ("F", "="): 0x0424,
+ ("H", "="): 0x0425,
+ ("C", "="): 0x0426,
+ ("C", "%"): 0x0427,
+ ("S", "%"): 0x0428,
+ ("S", "c"): 0x0429,
+ ("=", '"'): 0x042A,
+ ("Y", "="): 0x042B,
+ ("%", '"'): 0x042C,
+ ("J", "E"): 0x042D,
+ ("J", "U"): 0x042E,
+ ("J", "A"): 0x042F,
+ ("a", "="): 0x0430,
+ ("b", "="): 0x0431,
+ ("v", "="): 0x0432,
+ ("g", "="): 0x0433,
+ ("d", "="): 0x0434,
+ ("e", "="): 0x0435,
+ ("z", "%"): 0x0436,
+ ("z", "="): 0x0437,
+ ("i", "="): 0x0438,
+ ("j", "="): 0x0439,
+ ("k", "="): 0x043A,
+ ("l", "="): 0x043B,
+ ("m", "="): 0x043C,
+ ("n", "="): 0x043D,
+ ("o", "="): 0x043E,
+ ("p", "="): 0x043F,
+ ("r", "="): 0x0440,
+ ("s", "="): 0x0441,
+ ("t", "="): 0x0442,
+ ("u", "="): 0x0443,
+ ("f", "="): 0x0444,
+ ("h", "="): 0x0445,
+ ("c", "="): 0x0446,
+ ("c", "%"): 0x0447,
+ ("s", "%"): 0x0448,
+ ("s", "c"): 0x0449,
+ ("=", "'"): 0x044A,
+ ("y", "="): 0x044B,
+ ("%", "'"): 0x044C,
+ ("j", "e"): 0x044D,
+ ("j", "u"): 0x044E,
+ ("j", "a"): 0x044F,
+ ("i", "o"): 0x0451,
+ ("d", "%"): 0x0452,
+ ("g", "%"): 0x0453,
+ ("i", "e"): 0x0454,
+ ("d", "s"): 0x0455,
+ ("i", "i"): 0x0456,
+ ("y", "i"): 0x0457,
+ ("j", "%"): 0x0458,
+ ("l", "j"): 0x0459,
+ ("n", "j"): 0x045A,
+ ("t", "s"): 0x045B,
+ ("k", "j"): 0x045C,
+ ("v", "%"): 0x045E,
+ ("d", "z"): 0x045F,
+ ("Y", "3"): 0x0462,
+ ("y", "3"): 0x0463,
+ ("O", "3"): 0x046A,
+ ("o", "3"): 0x046B,
+ ("F", "3"): 0x0472,
+ ("f", "3"): 0x0473,
+ ("V", "3"): 0x0474,
+ ("v", "3"): 0x0475,
+ ("C", "3"): 0x0480,
+ ("c", "3"): 0x0481,
+ ("G", "3"): 0x0490,
+ ("g", "3"): 0x0491,
+ ("A", "+"): 0x05D0,
+ ("B", "+"): 0x05D1,
+ ("G", "+"): 0x05D2,
+ ("D", "+"): 0x05D3,
+ ("H", "+"): 0x05D4,
+ ("W", "+"): 0x05D5,
+ ("Z", "+"): 0x05D6,
+ ("X", "+"): 0x05D7,
+ ("T", "j"): 0x05D8,
+ ("J", "+"): 0x05D9,
+ ("K", "%"): 0x05DA,
+ ("K", "+"): 0x05DB,
+ ("L", "+"): 0x05DC,
+ ("M", "%"): 0x05DD,
+ ("M", "+"): 0x05DE,
+ ("N", "%"): 0x05DF,
+ ("N", "+"): 0x05E0,
+ ("S", "+"): 0x05E1,
+ ("E", "+"): 0x05E2,
+ ("P", "%"): 0x05E3,
+ ("P", "+"): 0x05E4,
+ ("Z", "j"): 0x05E5,
+ ("Z", "J"): 0x05E6,
+ ("Q", "+"): 0x05E7,
+ ("R", "+"): 0x05E8,
+ ("S", "h"): 0x05E9,
+ ("T", "+"): 0x05EA,
+ (",", "+"): 0x060C,
+ (";", "+"): 0x061B,
+ ("?", "+"): 0x061F,
+ ("H", "'"): 0x0621,
+ ("a", "M"): 0x0622,
+ ("a", "H"): 0x0623,
+ ("w", "H"): 0x0624,
+ ("a", "h"): 0x0625,
+ ("y", "H"): 0x0626,
+ ("a", "+"): 0x0627,
+ ("b", "+"): 0x0628,
+ ("t", "m"): 0x0629,
+ ("t", "+"): 0x062A,
+ ("t", "k"): 0x062B,
+ ("g", "+"): 0x062C,
+ ("h", "k"): 0x062D,
+ ("x", "+"): 0x062E,
+ ("d", "+"): 0x062F,
+ ("d", "k"): 0x0630,
+ ("r", "+"): 0x0631,
+ ("z", "+"): 0x0632,
+ ("s", "+"): 0x0633,
+ ("s", "n"): 0x0634,
+ ("c", "+"): 0x0635,
+ ("d", "d"): 0x0636,
+ ("t", "j"): 0x0637,
+ ("z", "H"): 0x0638,
+ ("e", "+"): 0x0639,
+ ("i", "+"): 0x063A,
+ ("+", "+"): 0x0640,
+ ("f", "+"): 0x0641,
+ ("q", "+"): 0x0642,
+ ("k", "+"): 0x0643,
+ ("l", "+"): 0x0644,
+ ("m", "+"): 0x0645,
+ ("n", "+"): 0x0646,
+ ("h", "+"): 0x0647,
+ ("w", "+"): 0x0648,
+ ("j", "+"): 0x0649,
+ ("y", "+"): 0x064A,
+ (":", "+"): 0x064B,
+ ('"', "+"): 0x064C,
+ ("=", "+"): 0x064D,
+ ("/", "+"): 0x064E,
+ ("'", "+"): 0x064F,
+ ("1", "+"): 0x0650,
+ ("3", "+"): 0x0651,
+ ("0", "+"): 0x0652,
+ ("a", "S"): 0x0670,
+ ("p", "+"): 0x067E,
+ ("v", "+"): 0x06A4,
+ ("g", "f"): 0x06AF,
+ ("0", "a"): 0x06F0,
+ ("1", "a"): 0x06F1,
+ ("2", "a"): 0x06F2,
+ ("3", "a"): 0x06F3,
+ ("4", "a"): 0x06F4,
+ ("5", "a"): 0x06F5,
+ ("6", "a"): 0x06F6,
+ ("7", "a"): 0x06F7,
+ ("8", "a"): 0x06F8,
+ ("9", "a"): 0x06F9,
+ ("B", "."): 0x1E02,
+ ("b", "."): 0x1E03,
+ ("B", "_"): 0x1E06,
+ ("b", "_"): 0x1E07,
+ ("D", "."): 0x1E0A,
+ ("d", "."): 0x1E0B,
+ ("D", "_"): 0x1E0E,
+ ("d", "_"): 0x1E0F,
+ ("D", ","): 0x1E10,
+ ("d", ","): 0x1E11,
+ ("F", "."): 0x1E1E,
+ ("f", "."): 0x1E1F,
+ ("G", "-"): 0x1E20,
+ ("g", "-"): 0x1E21,
+ ("H", "."): 0x1E22,
+ ("h", "."): 0x1E23,
+ ("H", ":"): 0x1E26,
+ ("h", ":"): 0x1E27,
+ ("H", ","): 0x1E28,
+ ("h", ","): 0x1E29,
+ ("K", "'"): 0x1E30,
+ ("k", "'"): 0x1E31,
+ ("K", "_"): 0x1E34,
+ ("k", "_"): 0x1E35,
+ ("L", "_"): 0x1E3A,
+ ("l", "_"): 0x1E3B,
+ ("M", "'"): 0x1E3E,
+ ("m", "'"): 0x1E3F,
+ ("M", "."): 0x1E40,
+ ("m", "."): 0x1E41,
+ ("N", "."): 0x1E44,
+ ("n", "."): 0x1E45,
+ ("N", "_"): 0x1E48,
+ ("n", "_"): 0x1E49,
+ ("P", "'"): 0x1E54,
+ ("p", "'"): 0x1E55,
+ ("P", "."): 0x1E56,
+ ("p", "."): 0x1E57,
+ ("R", "."): 0x1E58,
+ ("r", "."): 0x1E59,
+ ("R", "_"): 0x1E5E,
+ ("r", "_"): 0x1E5F,
+ ("S", "."): 0x1E60,
+ ("s", "."): 0x1E61,
+ ("T", "."): 0x1E6A,
+ ("t", "."): 0x1E6B,
+ ("T", "_"): 0x1E6E,
+ ("t", "_"): 0x1E6F,
+ ("V", "?"): 0x1E7C,
+ ("v", "?"): 0x1E7D,
+ ("W", "!"): 0x1E80,
+ ("w", "!"): 0x1E81,
+ ("W", "'"): 0x1E82,
+ ("w", "'"): 0x1E83,
+ ("W", ":"): 0x1E84,
+ ("w", ":"): 0x1E85,
+ ("W", "."): 0x1E86,
+ ("w", "."): 0x1E87,
+ ("X", "."): 0x1E8A,
+ ("x", "."): 0x1E8B,
+ ("X", ":"): 0x1E8C,
+ ("x", ":"): 0x1E8D,
+ ("Y", "."): 0x1E8E,
+ ("y", "."): 0x1E8F,
+ ("Z", ">"): 0x1E90,
+ ("z", ">"): 0x1E91,
+ ("Z", "_"): 0x1E94,
+ ("z", "_"): 0x1E95,
+ ("h", "_"): 0x1E96,
+ ("t", ":"): 0x1E97,
+ ("w", "0"): 0x1E98,
+ ("y", "0"): 0x1E99,
+ ("A", "2"): 0x1EA2,
+ ("a", "2"): 0x1EA3,
+ ("E", "2"): 0x1EBA,
+ ("e", "2"): 0x1EBB,
+ ("E", "?"): 0x1EBC,
+ ("e", "?"): 0x1EBD,
+ ("I", "2"): 0x1EC8,
+ ("i", "2"): 0x1EC9,
+ ("O", "2"): 0x1ECE,
+ ("o", "2"): 0x1ECF,
+ ("U", "2"): 0x1EE6,
+ ("u", "2"): 0x1EE7,
+ ("Y", "!"): 0x1EF2,
+ ("y", "!"): 0x1EF3,
+ ("Y", "2"): 0x1EF6,
+ ("y", "2"): 0x1EF7,
+ ("Y", "?"): 0x1EF8,
+ ("y", "?"): 0x1EF9,
+ (";", "'"): 0x1F00,
+ (",", "'"): 0x1F01,
+ (";", "!"): 0x1F02,
+ (",", "!"): 0x1F03,
+ ("?", ";"): 0x1F04,
+ ("?", ","): 0x1F05,
+ ("!", ":"): 0x1F06,
+ ("?", ":"): 0x1F07,
+ ("1", "N"): 0x2002,
+ ("1", "M"): 0x2003,
+ ("3", "M"): 0x2004,
+ ("4", "M"): 0x2005,
+ ("6", "M"): 0x2006,
+ ("1", "T"): 0x2009,
+ ("1", "H"): 0x200A,
+ ("-", "1"): 0x2010,
+ ("-", "N"): 0x2013,
+ ("-", "M"): 0x2014,
+ ("-", "3"): 0x2015,
+ ("!", "2"): 0x2016,
+ ("=", "2"): 0x2017,
+ ("'", "6"): 0x2018,
+ ("'", "9"): 0x2019,
+ (".", "9"): 0x201A,
+ ("9", "'"): 0x201B,
+ ('"', "6"): 0x201C,
+ ('"', "9"): 0x201D,
+ (":", "9"): 0x201E,
+ ("9", '"'): 0x201F,
+ ("/", "-"): 0x2020,
+ ("/", "="): 0x2021,
+ (".", "."): 0x2025,
+ ("%", "0"): 0x2030,
+ ("1", "'"): 0x2032,
+ ("2", "'"): 0x2033,
+ ("3", "'"): 0x2034,
+ ("1", '"'): 0x2035,
+ ("2", '"'): 0x2036,
+ ("3", '"'): 0x2037,
+ ("C", "a"): 0x2038,
+ ("<", "1"): 0x2039,
+ (">", "1"): 0x203A,
+ (":", "X"): 0x203B,
+ ("'", "-"): 0x203E,
+ ("/", "f"): 0x2044,
+ ("0", "S"): 0x2070,
+ ("4", "S"): 0x2074,
+ ("5", "S"): 0x2075,
+ ("6", "S"): 0x2076,
+ ("7", "S"): 0x2077,
+ ("8", "S"): 0x2078,
+ ("9", "S"): 0x2079,
+ ("+", "S"): 0x207A,
+ ("-", "S"): 0x207B,
+ ("=", "S"): 0x207C,
+ ("(", "S"): 0x207D,
+ (")", "S"): 0x207E,
+ ("n", "S"): 0x207F,
+ ("0", "s"): 0x2080,
+ ("1", "s"): 0x2081,
+ ("2", "s"): 0x2082,
+ ("3", "s"): 0x2083,
+ ("4", "s"): 0x2084,
+ ("5", "s"): 0x2085,
+ ("6", "s"): 0x2086,
+ ("7", "s"): 0x2087,
+ ("8", "s"): 0x2088,
+ ("9", "s"): 0x2089,
+ ("+", "s"): 0x208A,
+ ("-", "s"): 0x208B,
+ ("=", "s"): 0x208C,
+ ("(", "s"): 0x208D,
+ (")", "s"): 0x208E,
+ ("L", "i"): 0x20A4,
+ ("P", "t"): 0x20A7,
+ ("W", "="): 0x20A9,
+ ("=", "e"): 0x20AC, # euro
+ ("E", "u"): 0x20AC, # euro
+ ("=", "R"): 0x20BD, # rouble
+ ("=", "P"): 0x20BD, # rouble
+ ("o", "C"): 0x2103,
+ ("c", "o"): 0x2105,
+ ("o", "F"): 0x2109,
+ ("N", "0"): 0x2116,
+ ("P", "O"): 0x2117,
+ ("R", "x"): 0x211E,
+ ("S", "M"): 0x2120,
+ ("T", "M"): 0x2122,
+ ("O", "m"): 0x2126,
+ ("A", "O"): 0x212B,
+ ("1", "3"): 0x2153,
+ ("2", "3"): 0x2154,
+ ("1", "5"): 0x2155,
+ ("2", "5"): 0x2156,
+ ("3", "5"): 0x2157,
+ ("4", "5"): 0x2158,
+ ("1", "6"): 0x2159,
+ ("5", "6"): 0x215A,
+ ("1", "8"): 0x215B,
+ ("3", "8"): 0x215C,
+ ("5", "8"): 0x215D,
+ ("7", "8"): 0x215E,
+ ("1", "R"): 0x2160,
+ ("2", "R"): 0x2161,
+ ("3", "R"): 0x2162,
+ ("4", "R"): 0x2163,
+ ("5", "R"): 0x2164,
+ ("6", "R"): 0x2165,
+ ("7", "R"): 0x2166,
+ ("8", "R"): 0x2167,
+ ("9", "R"): 0x2168,
+ ("a", "R"): 0x2169,
+ ("b", "R"): 0x216A,
+ ("c", "R"): 0x216B,
+ ("1", "r"): 0x2170,
+ ("2", "r"): 0x2171,
+ ("3", "r"): 0x2172,
+ ("4", "r"): 0x2173,
+ ("5", "r"): 0x2174,
+ ("6", "r"): 0x2175,
+ ("7", "r"): 0x2176,
+ ("8", "r"): 0x2177,
+ ("9", "r"): 0x2178,
+ ("a", "r"): 0x2179,
+ ("b", "r"): 0x217A,
+ ("c", "r"): 0x217B,
+ ("<", "-"): 0x2190,
+ ("-", "!"): 0x2191,
+ ("-", ">"): 0x2192,
+ ("-", "v"): 0x2193,
+ ("<", ">"): 0x2194,
+ ("U", "D"): 0x2195,
+ ("<", "="): 0x21D0,
+ ("=", ">"): 0x21D2,
+ ("=", "="): 0x21D4,
+ ("F", "A"): 0x2200,
+ ("d", "P"): 0x2202,
+ ("T", "E"): 0x2203,
+ ("/", "0"): 0x2205,
+ ("D", "E"): 0x2206,
+ ("N", "B"): 0x2207,
+ ("(", "-"): 0x2208,
+ ("-", ")"): 0x220B,
+ ("*", "P"): 0x220F,
+ ("+", "Z"): 0x2211,
+ ("-", "2"): 0x2212,
+ ("-", "+"): 0x2213,
+ ("*", "-"): 0x2217,
+ ("O", "b"): 0x2218,
+ ("S", "b"): 0x2219,
+ ("R", "T"): 0x221A,
+ ("0", "("): 0x221D,
+ ("0", "0"): 0x221E,
+ ("-", "L"): 0x221F,
+ ("-", "V"): 0x2220,
+ ("P", "P"): 0x2225,
+ ("A", "N"): 0x2227,
+ ("O", "R"): 0x2228,
+ ("(", "U"): 0x2229,
+ (")", "U"): 0x222A,
+ ("I", "n"): 0x222B,
+ ("D", "I"): 0x222C,
+ ("I", "o"): 0x222E,
+ (".", ":"): 0x2234,
+ (":", "."): 0x2235,
+ (":", "R"): 0x2236,
+ (":", ":"): 0x2237,
+ ("?", "1"): 0x223C,
+ ("C", "G"): 0x223E,
+ ("?", "-"): 0x2243,
+ ("?", "="): 0x2245,
+ ("?", "2"): 0x2248,
+ ("=", "?"): 0x224C,
+ ("H", "I"): 0x2253,
+ ("!", "="): 0x2260,
+ ("=", "3"): 0x2261,
+ ("=", "<"): 0x2264,
+ (">", "="): 0x2265,
+ ("<", "*"): 0x226A,
+ ("*", ">"): 0x226B,
+ ("!", "<"): 0x226E,
+ ("!", ">"): 0x226F,
+ ("(", "C"): 0x2282,
+ (")", "C"): 0x2283,
+ ("(", "_"): 0x2286,
+ (")", "_"): 0x2287,
+ ("0", "."): 0x2299,
+ ("0", "2"): 0x229A,
+ ("-", "T"): 0x22A5,
+ (".", "P"): 0x22C5,
+ (":", "3"): 0x22EE,
+ (".", "3"): 0x22EF,
+ ("E", "h"): 0x2302,
+ ("<", "7"): 0x2308,
+ (">", "7"): 0x2309,
+ ("7", "<"): 0x230A,
+ ("7", ">"): 0x230B,
+ ("N", "I"): 0x2310,
+ ("(", "A"): 0x2312,
+ ("T", "R"): 0x2315,
+ ("I", "u"): 0x2320,
+ ("I", "l"): 0x2321,
+ ("<", "/"): 0x2329,
+ ("/", ">"): 0x232A,
+ ("V", "s"): 0x2423,
+ ("1", "h"): 0x2440,
+ ("3", "h"): 0x2441,
+ ("2", "h"): 0x2442,
+ ("4", "h"): 0x2443,
+ ("1", "j"): 0x2446,
+ ("2", "j"): 0x2447,
+ ("3", "j"): 0x2448,
+ ("4", "j"): 0x2449,
+ ("1", "."): 0x2488,
+ ("2", "."): 0x2489,
+ ("3", "."): 0x248A,
+ ("4", "."): 0x248B,
+ ("5", "."): 0x248C,
+ ("6", "."): 0x248D,
+ ("7", "."): 0x248E,
+ ("8", "."): 0x248F,
+ ("9", "."): 0x2490,
+ ("h", "h"): 0x2500,
+ ("H", "H"): 0x2501,
+ ("v", "v"): 0x2502,
+ ("V", "V"): 0x2503,
+ ("3", "-"): 0x2504,
+ ("3", "_"): 0x2505,
+ ("3", "!"): 0x2506,
+ ("3", "/"): 0x2507,
+ ("4", "-"): 0x2508,
+ ("4", "_"): 0x2509,
+ ("4", "!"): 0x250A,
+ ("4", "/"): 0x250B,
+ ("d", "r"): 0x250C,
+ ("d", "R"): 0x250D,
+ ("D", "r"): 0x250E,
+ ("D", "R"): 0x250F,
+ ("d", "l"): 0x2510,
+ ("d", "L"): 0x2511,
+ ("D", "l"): 0x2512,
+ ("L", "D"): 0x2513,
+ ("u", "r"): 0x2514,
+ ("u", "R"): 0x2515,
+ ("U", "r"): 0x2516,
+ ("U", "R"): 0x2517,
+ ("u", "l"): 0x2518,
+ ("u", "L"): 0x2519,
+ ("U", "l"): 0x251A,
+ ("U", "L"): 0x251B,
+ ("v", "r"): 0x251C,
+ ("v", "R"): 0x251D,
+ ("V", "r"): 0x2520,
+ ("V", "R"): 0x2523,
+ ("v", "l"): 0x2524,
+ ("v", "L"): 0x2525,
+ ("V", "l"): 0x2528,
+ ("V", "L"): 0x252B,
+ ("d", "h"): 0x252C,
+ ("d", "H"): 0x252F,
+ ("D", "h"): 0x2530,
+ ("D", "H"): 0x2533,
+ ("u", "h"): 0x2534,
+ ("u", "H"): 0x2537,
+ ("U", "h"): 0x2538,
+ ("U", "H"): 0x253B,
+ ("v", "h"): 0x253C,
+ ("v", "H"): 0x253F,
+ ("V", "h"): 0x2542,
+ ("V", "H"): 0x254B,
+ ("F", "D"): 0x2571,
+ ("B", "D"): 0x2572,
+ ("T", "B"): 0x2580,
+ ("L", "B"): 0x2584,
+ ("F", "B"): 0x2588,
+ ("l", "B"): 0x258C,
+ ("R", "B"): 0x2590,
+ (".", "S"): 0x2591,
+ (":", "S"): 0x2592,
+ ("?", "S"): 0x2593,
+ ("f", "S"): 0x25A0,
+ ("O", "S"): 0x25A1,
+ ("R", "O"): 0x25A2,
+ ("R", "r"): 0x25A3,
+ ("R", "F"): 0x25A4,
+ ("R", "Y"): 0x25A5,
+ ("R", "H"): 0x25A6,
+ ("R", "Z"): 0x25A7,
+ ("R", "K"): 0x25A8,
+ ("R", "X"): 0x25A9,
+ ("s", "B"): 0x25AA,
+ ("S", "R"): 0x25AC,
+ ("O", "r"): 0x25AD,
+ ("U", "T"): 0x25B2,
+ ("u", "T"): 0x25B3,
+ ("P", "R"): 0x25B6,
+ ("T", "r"): 0x25B7,
+ ("D", "t"): 0x25BC,
+ ("d", "T"): 0x25BD,
+ ("P", "L"): 0x25C0,
+ ("T", "l"): 0x25C1,
+ ("D", "b"): 0x25C6,
+ ("D", "w"): 0x25C7,
+ ("L", "Z"): 0x25CA,
+ ("0", "m"): 0x25CB,
+ ("0", "o"): 0x25CE,
+ ("0", "M"): 0x25CF,
+ ("0", "L"): 0x25D0,
+ ("0", "R"): 0x25D1,
+ ("S", "n"): 0x25D8,
+ ("I", "c"): 0x25D9,
+ ("F", "d"): 0x25E2,
+ ("B", "d"): 0x25E3,
+ ("*", "2"): 0x2605,
+ ("*", "1"): 0x2606,
+ ("<", "H"): 0x261C,
+ (">", "H"): 0x261E,
+ ("0", "u"): 0x263A,
+ ("0", "U"): 0x263B,
+ ("S", "U"): 0x263C,
+ ("F", "m"): 0x2640,
+ ("M", "l"): 0x2642,
+ ("c", "S"): 0x2660,
+ ("c", "H"): 0x2661,
+ ("c", "D"): 0x2662,
+ ("c", "C"): 0x2663,
+ ("M", "d"): 0x2669,
+ ("M", "8"): 0x266A,
+ ("M", "2"): 0x266B,
+ ("M", "b"): 0x266D,
+ ("M", "x"): 0x266E,
+ ("M", "X"): 0x266F,
+ ("O", "K"): 0x2713,
+ ("X", "X"): 0x2717,
+ ("-", "X"): 0x2720,
+ ("I", "S"): 0x3000,
+ (",", "_"): 0x3001,
+ (".", "_"): 0x3002,
+ ("+", '"'): 0x3003,
+ ("+", "_"): 0x3004,
+ ("*", "_"): 0x3005,
+ (";", "_"): 0x3006,
+ ("0", "_"): 0x3007,
+ ("<", "+"): 0x300A,
+ (">", "+"): 0x300B,
+ ("<", "'"): 0x300C,
+ (">", "'"): 0x300D,
+ ("<", '"'): 0x300E,
+ (">", '"'): 0x300F,
+ ("(", '"'): 0x3010,
+ (")", '"'): 0x3011,
+ ("=", "T"): 0x3012,
+ ("=", "_"): 0x3013,
+ ("(", "'"): 0x3014,
+ (")", "'"): 0x3015,
+ ("(", "I"): 0x3016,
+ (")", "I"): 0x3017,
+ ("-", "?"): 0x301C,
+ ("A", "5"): 0x3041,
+ ("a", "5"): 0x3042,
+ ("I", "5"): 0x3043,
+ ("i", "5"): 0x3044,
+ ("U", "5"): 0x3045,
+ ("u", "5"): 0x3046,
+ ("E", "5"): 0x3047,
+ ("e", "5"): 0x3048,
+ ("O", "5"): 0x3049,
+ ("o", "5"): 0x304A,
+ ("k", "a"): 0x304B,
+ ("g", "a"): 0x304C,
+ ("k", "i"): 0x304D,
+ ("g", "i"): 0x304E,
+ ("k", "u"): 0x304F,
+ ("g", "u"): 0x3050,
+ ("k", "e"): 0x3051,
+ ("g", "e"): 0x3052,
+ ("k", "o"): 0x3053,
+ ("g", "o"): 0x3054,
+ ("s", "a"): 0x3055,
+ ("z", "a"): 0x3056,
+ ("s", "i"): 0x3057,
+ ("z", "i"): 0x3058,
+ ("s", "u"): 0x3059,
+ ("z", "u"): 0x305A,
+ ("s", "e"): 0x305B,
+ ("z", "e"): 0x305C,
+ ("s", "o"): 0x305D,
+ ("z", "o"): 0x305E,
+ ("t", "a"): 0x305F,
+ ("d", "a"): 0x3060,
+ ("t", "i"): 0x3061,
+ ("d", "i"): 0x3062,
+ ("t", "U"): 0x3063,
+ ("t", "u"): 0x3064,
+ ("d", "u"): 0x3065,
+ ("t", "e"): 0x3066,
+ ("d", "e"): 0x3067,
+ ("t", "o"): 0x3068,
+ ("d", "o"): 0x3069,
+ ("n", "a"): 0x306A,
+ ("n", "i"): 0x306B,
+ ("n", "u"): 0x306C,
+ ("n", "e"): 0x306D,
+ ("n", "o"): 0x306E,
+ ("h", "a"): 0x306F,
+ ("b", "a"): 0x3070,
+ ("p", "a"): 0x3071,
+ ("h", "i"): 0x3072,
+ ("b", "i"): 0x3073,
+ ("p", "i"): 0x3074,
+ ("h", "u"): 0x3075,
+ ("b", "u"): 0x3076,
+ ("p", "u"): 0x3077,
+ ("h", "e"): 0x3078,
+ ("b", "e"): 0x3079,
+ ("p", "e"): 0x307A,
+ ("h", "o"): 0x307B,
+ ("b", "o"): 0x307C,
+ ("p", "o"): 0x307D,
+ ("m", "a"): 0x307E,
+ ("m", "i"): 0x307F,
+ ("m", "u"): 0x3080,
+ ("m", "e"): 0x3081,
+ ("m", "o"): 0x3082,
+ ("y", "A"): 0x3083,
+ ("y", "a"): 0x3084,
+ ("y", "U"): 0x3085,
+ ("y", "u"): 0x3086,
+ ("y", "O"): 0x3087,
+ ("y", "o"): 0x3088,
+ ("r", "a"): 0x3089,
+ ("r", "i"): 0x308A,
+ ("r", "u"): 0x308B,
+ ("r", "e"): 0x308C,
+ ("r", "o"): 0x308D,
+ ("w", "A"): 0x308E,
+ ("w", "a"): 0x308F,
+ ("w", "i"): 0x3090,
+ ("w", "e"): 0x3091,
+ ("w", "o"): 0x3092,
+ ("n", "5"): 0x3093,
+ ("v", "u"): 0x3094,
+ ('"', "5"): 0x309B,
+ ("0", "5"): 0x309C,
+ ("*", "5"): 0x309D,
+ ("+", "5"): 0x309E,
+ ("a", "6"): 0x30A1,
+ ("A", "6"): 0x30A2,
+ ("i", "6"): 0x30A3,
+ ("I", "6"): 0x30A4,
+ ("u", "6"): 0x30A5,
+ ("U", "6"): 0x30A6,
+ ("e", "6"): 0x30A7,
+ ("E", "6"): 0x30A8,
+ ("o", "6"): 0x30A9,
+ ("O", "6"): 0x30AA,
+ ("K", "a"): 0x30AB,
+ ("G", "a"): 0x30AC,
+ ("K", "i"): 0x30AD,
+ ("G", "i"): 0x30AE,
+ ("K", "u"): 0x30AF,
+ ("G", "u"): 0x30B0,
+ ("K", "e"): 0x30B1,
+ ("G", "e"): 0x30B2,
+ ("K", "o"): 0x30B3,
+ ("G", "o"): 0x30B4,
+ ("S", "a"): 0x30B5,
+ ("Z", "a"): 0x30B6,
+ ("S", "i"): 0x30B7,
+ ("Z", "i"): 0x30B8,
+ ("S", "u"): 0x30B9,
+ ("Z", "u"): 0x30BA,
+ ("S", "e"): 0x30BB,
+ ("Z", "e"): 0x30BC,
+ ("S", "o"): 0x30BD,
+ ("Z", "o"): 0x30BE,
+ ("T", "a"): 0x30BF,
+ ("D", "a"): 0x30C0,
+ ("T", "i"): 0x30C1,
+ ("D", "i"): 0x30C2,
+ ("T", "U"): 0x30C3,
+ ("T", "u"): 0x30C4,
+ ("D", "u"): 0x30C5,
+ ("T", "e"): 0x30C6,
+ ("D", "e"): 0x30C7,
+ ("T", "o"): 0x30C8,
+ ("D", "o"): 0x30C9,
+ ("N", "a"): 0x30CA,
+ ("N", "i"): 0x30CB,
+ ("N", "u"): 0x30CC,
+ ("N", "e"): 0x30CD,
+ ("N", "o"): 0x30CE,
+ ("H", "a"): 0x30CF,
+ ("B", "a"): 0x30D0,
+ ("P", "a"): 0x30D1,
+ ("H", "i"): 0x30D2,
+ ("B", "i"): 0x30D3,
+ ("P", "i"): 0x30D4,
+ ("H", "u"): 0x30D5,
+ ("B", "u"): 0x30D6,
+ ("P", "u"): 0x30D7,
+ ("H", "e"): 0x30D8,
+ ("B", "e"): 0x30D9,
+ ("P", "e"): 0x30DA,
+ ("H", "o"): 0x30DB,
+ ("B", "o"): 0x30DC,
+ ("P", "o"): 0x30DD,
+ ("M", "a"): 0x30DE,
+ ("M", "i"): 0x30DF,
+ ("M", "u"): 0x30E0,
+ ("M", "e"): 0x30E1,
+ ("M", "o"): 0x30E2,
+ ("Y", "A"): 0x30E3,
+ ("Y", "a"): 0x30E4,
+ ("Y", "U"): 0x30E5,
+ ("Y", "u"): 0x30E6,
+ ("Y", "O"): 0x30E7,
+ ("Y", "o"): 0x30E8,
+ ("R", "a"): 0x30E9,
+ ("R", "i"): 0x30EA,
+ ("R", "u"): 0x30EB,
+ ("R", "e"): 0x30EC,
+ ("R", "o"): 0x30ED,
+ ("W", "A"): 0x30EE,
+ ("W", "a"): 0x30EF,
+ ("W", "i"): 0x30F0,
+ ("W", "e"): 0x30F1,
+ ("W", "o"): 0x30F2,
+ ("N", "6"): 0x30F3,
+ ("V", "u"): 0x30F4,
+ ("K", "A"): 0x30F5,
+ ("K", "E"): 0x30F6,
+ ("V", "a"): 0x30F7,
+ ("V", "i"): 0x30F8,
+ ("V", "e"): 0x30F9,
+ ("V", "o"): 0x30FA,
+ (".", "6"): 0x30FB,
+ ("-", "6"): 0x30FC,
+ ("*", "6"): 0x30FD,
+ ("+", "6"): 0x30FE,
+ ("b", "4"): 0x3105,
+ ("p", "4"): 0x3106,
+ ("m", "4"): 0x3107,
+ ("f", "4"): 0x3108,
+ ("d", "4"): 0x3109,
+ ("t", "4"): 0x310A,
+ ("n", "4"): 0x310B,
+ ("l", "4"): 0x310C,
+ ("g", "4"): 0x310D,
+ ("k", "4"): 0x310E,
+ ("h", "4"): 0x310F,
+ ("j", "4"): 0x3110,
+ ("q", "4"): 0x3111,
+ ("x", "4"): 0x3112,
+ ("z", "h"): 0x3113,
+ ("c", "h"): 0x3114,
+ ("s", "h"): 0x3115,
+ ("r", "4"): 0x3116,
+ ("z", "4"): 0x3117,
+ ("c", "4"): 0x3118,
+ ("s", "4"): 0x3119,
+ ("a", "4"): 0x311A,
+ ("o", "4"): 0x311B,
+ ("e", "4"): 0x311C,
+ ("a", "i"): 0x311E,
+ ("e", "i"): 0x311F,
+ ("a", "u"): 0x3120,
+ ("o", "u"): 0x3121,
+ ("a", "n"): 0x3122,
+ ("e", "n"): 0x3123,
+ ("a", "N"): 0x3124,
+ ("e", "N"): 0x3125,
+ ("e", "r"): 0x3126,
+ ("i", "4"): 0x3127,
+ ("u", "4"): 0x3128,
+ ("i", "u"): 0x3129,
+ ("v", "4"): 0x312A,
+ ("n", "G"): 0x312B,
+ ("g", "n"): 0x312C,
+ ("1", "c"): 0x3220,
+ ("2", "c"): 0x3221,
+ ("3", "c"): 0x3222,
+ ("4", "c"): 0x3223,
+ ("5", "c"): 0x3224,
+ ("6", "c"): 0x3225,
+ ("7", "c"): 0x3226,
+ ("8", "c"): 0x3227,
+ ("9", "c"): 0x3228,
+ # code points 0xe000 - 0xefff excluded, they have no assigned
+ # characters, only used in proposals.
+ ("f", "f"): 0xFB00,
+ ("f", "i"): 0xFB01,
+ ("f", "l"): 0xFB02,
+ ("f", "t"): 0xFB05,
+ ("s", "t"): 0xFB06,
+ # Vim 5.x compatible digraphs that don't conflict with the above
+ ("~", "!"): 161,
+ ("c", "|"): 162,
+ ("$", "$"): 163,
+ ("o", "x"): 164, # currency symbol in ISO 8859-1
+ ("Y", "-"): 165,
+ ("|", "|"): 166,
+ ("c", "O"): 169,
+ ("-", ","): 172,
+ ("-", "="): 175,
+ ("~", "o"): 176,
+ ("2", "2"): 178,
+ ("3", "3"): 179,
+ ("p", "p"): 182,
+ ("~", "."): 183,
+ ("1", "1"): 185,
+ ("~", "?"): 191,
+ ("A", "`"): 192,
+ ("A", "^"): 194,
+ ("A", "~"): 195,
+ ("A", '"'): 196,
+ ("A", "@"): 197,
+ ("E", "`"): 200,
+ ("E", "^"): 202,
+ ("E", '"'): 203,
+ ("I", "`"): 204,
+ ("I", "^"): 206,
+ ("I", '"'): 207,
+ ("N", "~"): 209,
+ ("O", "`"): 210,
+ ("O", "^"): 212,
+ ("O", "~"): 213,
+ ("/", "\\"): 215, # multiplication symbol in ISO 8859-1
+ ("U", "`"): 217,
+ ("U", "^"): 219,
+ ("I", "p"): 222,
+ ("a", "`"): 224,
+ ("a", "^"): 226,
+ ("a", "~"): 227,
+ ("a", '"'): 228,
+ ("a", "@"): 229,
+ ("e", "`"): 232,
+ ("e", "^"): 234,
+ ("e", '"'): 235,
+ ("i", "`"): 236,
+ ("i", "^"): 238,
+ ("n", "~"): 241,
+ ("o", "`"): 242,
+ ("o", "^"): 244,
+ ("o", "~"): 245,
+ ("u", "`"): 249,
+ ("u", "^"): 251,
+ ("y", '"'): 255,
+}
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/emacs_state.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/emacs_state.py
new file mode 100644
index 0000000000000000000000000000000000000000..6a2ebf46d803e03bfab689a27c6e48f437ab355e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/emacs_state.py
@@ -0,0 +1,36 @@
+from __future__ import annotations
+
+from .key_processor import KeyPress
+
+__all__ = [
+ "EmacsState",
+]
+
+
+class EmacsState:
+ """
+ Mutable class to hold Emacs specific state.
+ """
+
+ def __init__(self) -> None:
+ # Simple macro recording. (Like Readline does.)
+ # (For Emacs mode.)
+ self.macro: list[KeyPress] | None = []
+ self.current_recording: list[KeyPress] | None = None
+
+ def reset(self) -> None:
+ self.current_recording = None
+
+ @property
+ def is_recording(self) -> bool:
+ "Tell whether we are recording a macro."
+ return self.current_recording is not None
+
+ def start_macro(self) -> None:
+ "Start recording macro."
+ self.current_recording = []
+
+ def end_macro(self) -> None:
+ "End recording macro."
+ self.macro = self.current_recording
+ self.current_recording = None
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/key_bindings.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/key_bindings.py
new file mode 100644
index 0000000000000000000000000000000000000000..cf37c6d90d6446692939e41a40a70ec373cb567e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/key_bindings.py
@@ -0,0 +1,672 @@
+"""
+Key bindings registry.
+
+A `KeyBindings` object is a container that holds a list of key bindings. It has a
+very efficient internal data structure for checking which key bindings apply
+for a pressed key.
+
+Typical usage::
+
+ kb = KeyBindings()
+
+ @kb.add(Keys.ControlX, Keys.ControlC, filter=INSERT)
+ def handler(event):
+ # Handle ControlX-ControlC key sequence.
+ pass
+
+It is also possible to combine multiple KeyBindings objects. We do this in the
+default key bindings. There are some KeyBindings objects that contain the Emacs
+bindings, while others contain the Vi bindings. They are merged together using
+`merge_key_bindings`.
+
+We also have a `ConditionalKeyBindings` object that can enable/disable a group of
+key bindings at once.
+
+
+It is also possible to add a filter to a function, before a key binding has
+been assigned, through the `key_binding` decorator.::
+
+ # First define a key handler with the `filter`.
+ @key_binding(filter=condition)
+ def my_key_binding(event):
+ ...
+
+ # Later, add it to the key bindings.
+ kb.add(Keys.A, my_key_binding)
+"""
+
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from inspect import isawaitable
+from typing import (
+ TYPE_CHECKING,
+ Any,
+ Callable,
+ Coroutine,
+ Hashable,
+ Sequence,
+ Tuple,
+ TypeVar,
+ Union,
+ cast,
+)
+
+from prompt_toolkit.cache import SimpleCache
+from prompt_toolkit.filters import FilterOrBool, Never, to_filter
+from prompt_toolkit.keys import KEY_ALIASES, Keys
+
+if TYPE_CHECKING:
+ # Avoid circular imports.
+ from .key_processor import KeyPressEvent
+
+ # The only two return values for a mouse handler (and key bindings) are
+ # `None` and `NotImplemented`. For the type checker it's best to annotate
+ # this as `object`. (The consumer never expects a more specific instance:
+ # checking for NotImplemented can be done using `is NotImplemented`.)
+ NotImplementedOrNone = object
+ # Other non-working options are:
+ # * Optional[Literal[NotImplemented]]
+ # --> Doesn't work, Literal can't take an Any.
+ # * None
+ # --> Doesn't work. We can't assign the result of a function that
+ # returns `None` to a variable.
+ # * Any
+ # --> Works, but too broad.
+
+
+__all__ = [
+ "NotImplementedOrNone",
+ "Binding",
+ "KeyBindingsBase",
+ "KeyBindings",
+ "ConditionalKeyBindings",
+ "merge_key_bindings",
+ "DynamicKeyBindings",
+ "GlobalOnlyKeyBindings",
+]
+
+# Key bindings can be regular functions or coroutines.
+# In both cases, if they return `NotImplemented`, the UI won't be invalidated.
+# This is mainly used in case of mouse move events, to prevent excessive
+# repainting during mouse move events.
+KeyHandlerCallable = Callable[
+ ["KeyPressEvent"],
+ Union["NotImplementedOrNone", Coroutine[Any, Any, "NotImplementedOrNone"]],
+]
+
+
+class Binding:
+ """
+ Key binding: (key sequence + handler + filter).
+ (Immutable binding class.)
+
+ :param record_in_macro: When True, don't record this key binding when a
+ macro is recorded.
+ """
+
+ def __init__(
+ self,
+ keys: tuple[Keys | str, ...],
+ handler: KeyHandlerCallable,
+ filter: FilterOrBool = True,
+ eager: FilterOrBool = False,
+ is_global: FilterOrBool = False,
+ save_before: Callable[[KeyPressEvent], bool] = (lambda e: True),
+ record_in_macro: FilterOrBool = True,
+ ) -> None:
+ self.keys = keys
+ self.handler = handler
+ self.filter = to_filter(filter)
+ self.eager = to_filter(eager)
+ self.is_global = to_filter(is_global)
+ self.save_before = save_before
+ self.record_in_macro = to_filter(record_in_macro)
+
+ def call(self, event: KeyPressEvent) -> None:
+ result = self.handler(event)
+
+ # If the handler is a coroutine, create an asyncio task.
+ if isawaitable(result):
+ awaitable = cast(Coroutine[Any, Any, "NotImplementedOrNone"], result)
+
+ async def bg_task() -> None:
+ result = await awaitable
+ if result != NotImplemented:
+ event.app.invalidate()
+
+ event.app.create_background_task(bg_task())
+
+ elif result != NotImplemented:
+ event.app.invalidate()
+
+ def __repr__(self) -> str:
+ return (
+ f"{self.__class__.__name__}(keys={self.keys!r}, handler={self.handler!r})"
+ )
+
+
+# Sequence of keys presses.
+KeysTuple = Tuple[Union[Keys, str], ...]
+
+
+class KeyBindingsBase(metaclass=ABCMeta):
+ """
+ Interface for a KeyBindings.
+ """
+
+ @property
+ @abstractmethod
+ def _version(self) -> Hashable:
+ """
+ For cache invalidation. - This should increase every time that
+ something changes.
+ """
+ return 0
+
+ @abstractmethod
+ def get_bindings_for_keys(self, keys: KeysTuple) -> list[Binding]:
+ """
+ Return a list of key bindings that can handle these keys.
+ (This return also inactive bindings, so the `filter` still has to be
+ called, for checking it.)
+
+ :param keys: tuple of keys.
+ """
+ return []
+
+ @abstractmethod
+ def get_bindings_starting_with_keys(self, keys: KeysTuple) -> list[Binding]:
+ """
+ Return a list of key bindings that handle a key sequence starting with
+ `keys`. (It does only return bindings for which the sequences are
+ longer than `keys`. And like `get_bindings_for_keys`, it also includes
+ inactive bindings.)
+
+ :param keys: tuple of keys.
+ """
+ return []
+
+ @property
+ @abstractmethod
+ def bindings(self) -> list[Binding]:
+ """
+ List of `Binding` objects.
+ (These need to be exposed, so that `KeyBindings` objects can be merged
+ together.)
+ """
+ return []
+
+ # `add` and `remove` don't have to be part of this interface.
+
+
+T = TypeVar("T", bound=Union[KeyHandlerCallable, Binding])
+
+
+class KeyBindings(KeyBindingsBase):
+ """
+ A container for a set of key bindings.
+
+ Example usage::
+
+ kb = KeyBindings()
+
+ @kb.add('c-t')
+ def _(event):
+ print('Control-T pressed')
+
+ @kb.add('c-a', 'c-b')
+ def _(event):
+ print('Control-A pressed, followed by Control-B')
+
+ @kb.add('c-x', filter=is_searching)
+ def _(event):
+ print('Control-X pressed') # Works only if we are searching.
+
+ """
+
+ def __init__(self) -> None:
+ self._bindings: list[Binding] = []
+ self._get_bindings_for_keys_cache: SimpleCache[KeysTuple, list[Binding]] = (
+ SimpleCache(maxsize=10000)
+ )
+ self._get_bindings_starting_with_keys_cache: SimpleCache[
+ KeysTuple, list[Binding]
+ ] = SimpleCache(maxsize=1000)
+ self.__version = 0 # For cache invalidation.
+
+ def _clear_cache(self) -> None:
+ self.__version += 1
+ self._get_bindings_for_keys_cache.clear()
+ self._get_bindings_starting_with_keys_cache.clear()
+
+ @property
+ def bindings(self) -> list[Binding]:
+ return self._bindings
+
+ @property
+ def _version(self) -> Hashable:
+ return self.__version
+
+ def add(
+ self,
+ *keys: Keys | str,
+ filter: FilterOrBool = True,
+ eager: FilterOrBool = False,
+ is_global: FilterOrBool = False,
+ save_before: Callable[[KeyPressEvent], bool] = (lambda e: True),
+ record_in_macro: FilterOrBool = True,
+ ) -> Callable[[T], T]:
+ """
+ Decorator for adding a key bindings.
+
+ :param filter: :class:`~prompt_toolkit.filters.Filter` to determine
+ when this key binding is active.
+ :param eager: :class:`~prompt_toolkit.filters.Filter` or `bool`.
+ When True, ignore potential longer matches when this key binding is
+ hit. E.g. when there is an active eager key binding for Ctrl-X,
+ execute the handler immediately and ignore the key binding for
+ Ctrl-X Ctrl-E of which it is a prefix.
+ :param is_global: When this key bindings is added to a `Container` or
+ `Control`, make it a global (always active) binding.
+ :param save_before: Callable that takes an `Event` and returns True if
+ we should save the current buffer, before handling the event.
+ (That's the default.)
+ :param record_in_macro: Record these key bindings when a macro is
+ being recorded. (True by default.)
+ """
+ assert keys
+
+ keys = tuple(_parse_key(k) for k in keys)
+
+ if isinstance(filter, Never):
+ # When a filter is Never, it will always stay disabled, so in that
+ # case don't bother putting it in the key bindings. It will slow
+ # down every key press otherwise.
+ def decorator(func: T) -> T:
+ return func
+
+ else:
+
+ def decorator(func: T) -> T:
+ if isinstance(func, Binding):
+ # We're adding an existing Binding object.
+ self.bindings.append(
+ Binding(
+ keys,
+ func.handler,
+ filter=func.filter & to_filter(filter),
+ eager=to_filter(eager) | func.eager,
+ is_global=to_filter(is_global) | func.is_global,
+ save_before=func.save_before,
+ record_in_macro=func.record_in_macro,
+ )
+ )
+ else:
+ self.bindings.append(
+ Binding(
+ keys,
+ cast(KeyHandlerCallable, func),
+ filter=filter,
+ eager=eager,
+ is_global=is_global,
+ save_before=save_before,
+ record_in_macro=record_in_macro,
+ )
+ )
+ self._clear_cache()
+
+ return func
+
+ return decorator
+
+ def remove(self, *args: Keys | str | KeyHandlerCallable) -> None:
+ """
+ Remove a key binding.
+
+ This expects either a function that was given to `add` method as
+ parameter or a sequence of key bindings.
+
+ Raises `ValueError` when no bindings was found.
+
+ Usage::
+
+ remove(handler) # Pass handler.
+ remove('c-x', 'c-a') # Or pass the key bindings.
+ """
+ found = False
+
+ if callable(args[0]):
+ assert len(args) == 1
+ function = args[0]
+
+ # Remove the given function.
+ for b in self.bindings:
+ if b.handler == function:
+ self.bindings.remove(b)
+ found = True
+
+ else:
+ assert len(args) > 0
+ args = cast(Tuple[Union[Keys, str]], args)
+
+ # Remove this sequence of key bindings.
+ keys = tuple(_parse_key(k) for k in args)
+
+ for b in self.bindings:
+ if b.keys == keys:
+ self.bindings.remove(b)
+ found = True
+
+ if found:
+ self._clear_cache()
+ else:
+ # No key binding found for this function. Raise ValueError.
+ raise ValueError(f"Binding not found: {function!r}")
+
+ # For backwards-compatibility.
+ add_binding = add
+ remove_binding = remove
+
+ def get_bindings_for_keys(self, keys: KeysTuple) -> list[Binding]:
+ """
+ Return a list of key bindings that can handle this key.
+ (This return also inactive bindings, so the `filter` still has to be
+ called, for checking it.)
+
+ :param keys: tuple of keys.
+ """
+
+ def get() -> list[Binding]:
+ result: list[tuple[int, Binding]] = []
+
+ for b in self.bindings:
+ if len(keys) == len(b.keys):
+ match = True
+ any_count = 0
+
+ for i, j in zip(b.keys, keys):
+ if i != j and i != Keys.Any:
+ match = False
+ break
+
+ if i == Keys.Any:
+ any_count += 1
+
+ if match:
+ result.append((any_count, b))
+
+ # Place bindings that have more 'Any' occurrences in them at the end.
+ result = sorted(result, key=lambda item: -item[0])
+
+ return [item[1] for item in result]
+
+ return self._get_bindings_for_keys_cache.get(keys, get)
+
+ def get_bindings_starting_with_keys(self, keys: KeysTuple) -> list[Binding]:
+ """
+ Return a list of key bindings that handle a key sequence starting with
+ `keys`. (It does only return bindings for which the sequences are
+ longer than `keys`. And like `get_bindings_for_keys`, it also includes
+ inactive bindings.)
+
+ :param keys: tuple of keys.
+ """
+
+ def get() -> list[Binding]:
+ result = []
+ for b in self.bindings:
+ if len(keys) < len(b.keys):
+ match = True
+ for i, j in zip(b.keys, keys):
+ if i != j and i != Keys.Any:
+ match = False
+ break
+ if match:
+ result.append(b)
+ return result
+
+ return self._get_bindings_starting_with_keys_cache.get(keys, get)
+
+
+def _parse_key(key: Keys | str) -> str | Keys:
+ """
+ Replace key by alias and verify whether it's a valid one.
+ """
+ # Already a parse key? -> Return it.
+ if isinstance(key, Keys):
+ return key
+
+ # Lookup aliases.
+ key = KEY_ALIASES.get(key, key)
+
+ # Replace 'space' by ' '
+ if key == "space":
+ key = " "
+
+ # Return as `Key` object when it's a special key.
+ try:
+ return Keys(key)
+ except ValueError:
+ pass
+
+ # Final validation.
+ if len(key) != 1:
+ raise ValueError(f"Invalid key: {key}")
+
+ return key
+
+
+def key_binding(
+ filter: FilterOrBool = True,
+ eager: FilterOrBool = False,
+ is_global: FilterOrBool = False,
+ save_before: Callable[[KeyPressEvent], bool] = (lambda event: True),
+ record_in_macro: FilterOrBool = True,
+) -> Callable[[KeyHandlerCallable], Binding]:
+ """
+ Decorator that turn a function into a `Binding` object. This can be added
+ to a `KeyBindings` object when a key binding is assigned.
+ """
+ assert save_before is None or callable(save_before)
+
+ filter = to_filter(filter)
+ eager = to_filter(eager)
+ is_global = to_filter(is_global)
+ save_before = save_before
+ record_in_macro = to_filter(record_in_macro)
+ keys = ()
+
+ def decorator(function: KeyHandlerCallable) -> Binding:
+ return Binding(
+ keys,
+ function,
+ filter=filter,
+ eager=eager,
+ is_global=is_global,
+ save_before=save_before,
+ record_in_macro=record_in_macro,
+ )
+
+ return decorator
+
+
+class _Proxy(KeyBindingsBase):
+ """
+ Common part for ConditionalKeyBindings and _MergedKeyBindings.
+ """
+
+ def __init__(self) -> None:
+ # `KeyBindings` to be synchronized with all the others.
+ self._bindings2: KeyBindingsBase = KeyBindings()
+ self._last_version: Hashable = ()
+
+ def _update_cache(self) -> None:
+ """
+ If `self._last_version` is outdated, then this should update
+ the version and `self._bindings2`.
+ """
+ raise NotImplementedError
+
+ # Proxy methods to self._bindings2.
+
+ @property
+ def bindings(self) -> list[Binding]:
+ self._update_cache()
+ return self._bindings2.bindings
+
+ @property
+ def _version(self) -> Hashable:
+ self._update_cache()
+ return self._last_version
+
+ def get_bindings_for_keys(self, keys: KeysTuple) -> list[Binding]:
+ self._update_cache()
+ return self._bindings2.get_bindings_for_keys(keys)
+
+ def get_bindings_starting_with_keys(self, keys: KeysTuple) -> list[Binding]:
+ self._update_cache()
+ return self._bindings2.get_bindings_starting_with_keys(keys)
+
+
+class ConditionalKeyBindings(_Proxy):
+ """
+ Wraps around a `KeyBindings`. Disable/enable all the key bindings according to
+ the given (additional) filter.::
+
+ @Condition
+ def setting_is_true():
+ return True # or False
+
+ registry = ConditionalKeyBindings(key_bindings, setting_is_true)
+
+ When new key bindings are added to this object. They are also
+ enable/disabled according to the given `filter`.
+
+ :param registries: List of :class:`.KeyBindings` objects.
+ :param filter: :class:`~prompt_toolkit.filters.Filter` object.
+ """
+
+ def __init__(
+ self, key_bindings: KeyBindingsBase, filter: FilterOrBool = True
+ ) -> None:
+ _Proxy.__init__(self)
+
+ self.key_bindings = key_bindings
+ self.filter = to_filter(filter)
+
+ def _update_cache(self) -> None:
+ "If the original key bindings was changed. Update our copy version."
+ expected_version = self.key_bindings._version
+
+ if self._last_version != expected_version:
+ bindings2 = KeyBindings()
+
+ # Copy all bindings from `self.key_bindings`, adding our condition.
+ for b in self.key_bindings.bindings:
+ bindings2.bindings.append(
+ Binding(
+ keys=b.keys,
+ handler=b.handler,
+ filter=self.filter & b.filter,
+ eager=b.eager,
+ is_global=b.is_global,
+ save_before=b.save_before,
+ record_in_macro=b.record_in_macro,
+ )
+ )
+
+ self._bindings2 = bindings2
+ self._last_version = expected_version
+
+
+class _MergedKeyBindings(_Proxy):
+ """
+ Merge multiple registries of key bindings into one.
+
+ This class acts as a proxy to multiple :class:`.KeyBindings` objects, but
+ behaves as if this is just one bigger :class:`.KeyBindings`.
+
+ :param registries: List of :class:`.KeyBindings` objects.
+ """
+
+ def __init__(self, registries: Sequence[KeyBindingsBase]) -> None:
+ _Proxy.__init__(self)
+ self.registries = registries
+
+ def _update_cache(self) -> None:
+ """
+ If one of the original registries was changed. Update our merged
+ version.
+ """
+ expected_version = tuple(r._version for r in self.registries)
+
+ if self._last_version != expected_version:
+ bindings2 = KeyBindings()
+
+ for reg in self.registries:
+ bindings2.bindings.extend(reg.bindings)
+
+ self._bindings2 = bindings2
+ self._last_version = expected_version
+
+
+def merge_key_bindings(bindings: Sequence[KeyBindingsBase]) -> _MergedKeyBindings:
+ """
+ Merge multiple :class:`.Keybinding` objects together.
+
+ Usage::
+
+ bindings = merge_key_bindings([bindings1, bindings2, ...])
+ """
+ return _MergedKeyBindings(bindings)
+
+
+class DynamicKeyBindings(_Proxy):
+ """
+ KeyBindings class that can dynamically returns any KeyBindings.
+
+ :param get_key_bindings: Callable that returns a :class:`.KeyBindings` instance.
+ """
+
+ def __init__(self, get_key_bindings: Callable[[], KeyBindingsBase | None]) -> None:
+ self.get_key_bindings = get_key_bindings
+ self.__version = 0
+ self._last_child_version = None
+ self._dummy = KeyBindings() # Empty key bindings.
+
+ def _update_cache(self) -> None:
+ key_bindings = self.get_key_bindings() or self._dummy
+ assert isinstance(key_bindings, KeyBindingsBase)
+ version = id(key_bindings), key_bindings._version
+
+ self._bindings2 = key_bindings
+ self._last_version = version
+
+
+class GlobalOnlyKeyBindings(_Proxy):
+ """
+ Wrapper around a :class:`.KeyBindings` object that only exposes the global
+ key bindings.
+ """
+
+ def __init__(self, key_bindings: KeyBindingsBase) -> None:
+ _Proxy.__init__(self)
+ self.key_bindings = key_bindings
+
+ def _update_cache(self) -> None:
+ """
+ If one of the original registries was changed. Update our merged
+ version.
+ """
+ expected_version = self.key_bindings._version
+
+ if self._last_version != expected_version:
+ bindings2 = KeyBindings()
+
+ for b in self.key_bindings.bindings:
+ if b.is_global():
+ bindings2.bindings.append(b)
+
+ self._bindings2 = bindings2
+ self._last_version = expected_version
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/key_processor.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/key_processor.py
new file mode 100644
index 0000000000000000000000000000000000000000..e2070a14f02b07be7c0ea8f5b88e3aa9898a5a4d
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/key_processor.py
@@ -0,0 +1,526 @@
+"""
+An :class:`~.KeyProcessor` receives callbacks for the keystrokes parsed from
+the input in the :class:`~prompt_toolkit.inputstream.InputStream` instance.
+
+The `KeyProcessor` will according to the implemented keybindings call the
+correct callbacks when new key presses are feed through `feed`.
+"""
+
+from __future__ import annotations
+
+import weakref
+from asyncio import Task, sleep
+from collections import deque
+from typing import TYPE_CHECKING, Any, Generator
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.enums import EditingMode
+from prompt_toolkit.filters.app import vi_navigation_mode
+from prompt_toolkit.keys import Keys
+from prompt_toolkit.utils import Event
+
+from .key_bindings import Binding, KeyBindingsBase
+
+if TYPE_CHECKING:
+ from prompt_toolkit.application import Application
+ from prompt_toolkit.buffer import Buffer
+
+
+__all__ = [
+ "KeyProcessor",
+ "KeyPress",
+ "KeyPressEvent",
+]
+
+
+class KeyPress:
+ """
+ :param key: A `Keys` instance or text (one character).
+ :param data: The received string on stdin. (Often vt100 escape codes.)
+ """
+
+ def __init__(self, key: Keys | str, data: str | None = None) -> None:
+ assert isinstance(key, Keys) or len(key) == 1
+
+ if data is None:
+ if isinstance(key, Keys):
+ data = key.value
+ else:
+ data = key # 'key' is a one character string.
+
+ self.key = key
+ self.data = data
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(key={self.key!r}, data={self.data!r})"
+
+ def __eq__(self, other: object) -> bool:
+ if not isinstance(other, KeyPress):
+ return False
+ return self.key == other.key and self.data == other.data
+
+
+"""
+Helper object to indicate flush operation in the KeyProcessor.
+NOTE: the implementation is very similar to the VT100 parser.
+"""
+_Flush = KeyPress("?", data="_Flush")
+
+
+class KeyProcessor:
+ """
+ Statemachine that receives :class:`KeyPress` instances and according to the
+ key bindings in the given :class:`KeyBindings`, calls the matching handlers.
+
+ ::
+
+ p = KeyProcessor(key_bindings)
+
+ # Send keys into the processor.
+ p.feed(KeyPress(Keys.ControlX, '\x18'))
+ p.feed(KeyPress(Keys.ControlC, '\x03')
+
+ # Process all the keys in the queue.
+ p.process_keys()
+
+ # Now the ControlX-ControlC callback will be called if this sequence is
+ # registered in the key bindings.
+
+ :param key_bindings: `KeyBindingsBase` instance.
+ """
+
+ def __init__(self, key_bindings: KeyBindingsBase) -> None:
+ self._bindings = key_bindings
+
+ self.before_key_press = Event(self)
+ self.after_key_press = Event(self)
+
+ self._flush_wait_task: Task[None] | None = None
+
+ self.reset()
+
+ def reset(self) -> None:
+ self._previous_key_sequence: list[KeyPress] = []
+ self._previous_handler: Binding | None = None
+
+ # The queue of keys not yet send to our _process generator/state machine.
+ self.input_queue: deque[KeyPress] = deque()
+
+ # The key buffer that is matched in the generator state machine.
+ # (This is at at most the amount of keys that make up for one key binding.)
+ self.key_buffer: list[KeyPress] = []
+
+ #: Readline argument (for repetition of commands.)
+ #: https://www.gnu.org/software/bash/manual/html_node/Readline-Arguments.html
+ self.arg: str | None = None
+
+ # Start the processor coroutine.
+ self._process_coroutine = self._process()
+ self._process_coroutine.send(None) # type: ignore
+
+ def _get_matches(self, key_presses: list[KeyPress]) -> list[Binding]:
+ """
+ For a list of :class:`KeyPress` instances. Give the matching handlers
+ that would handle this.
+ """
+ keys = tuple(k.key for k in key_presses)
+
+ # Try match, with mode flag
+ return [b for b in self._bindings.get_bindings_for_keys(keys) if b.filter()]
+
+ def _is_prefix_of_longer_match(self, key_presses: list[KeyPress]) -> bool:
+ """
+ For a list of :class:`KeyPress` instances. Return True if there is any
+ handler that is bound to a suffix of this keys.
+ """
+ keys = tuple(k.key for k in key_presses)
+
+ # Get the filters for all the key bindings that have a longer match.
+ # Note that we transform it into a `set`, because we don't care about
+ # the actual bindings and executing it more than once doesn't make
+ # sense. (Many key bindings share the same filter.)
+ filters = {
+ b.filter for b in self._bindings.get_bindings_starting_with_keys(keys)
+ }
+
+ # When any key binding is active, return True.
+ return any(f() for f in filters)
+
+ def _process(self) -> Generator[None, KeyPress, None]:
+ """
+ Coroutine implementing the key match algorithm. Key strokes are sent
+ into this generator, and it calls the appropriate handlers.
+ """
+ buffer = self.key_buffer
+ retry = False
+
+ while True:
+ flush = False
+
+ if retry:
+ retry = False
+ else:
+ key = yield
+ if key is _Flush:
+ flush = True
+ else:
+ buffer.append(key)
+
+ # If we have some key presses, check for matches.
+ if buffer:
+ matches = self._get_matches(buffer)
+
+ if flush:
+ is_prefix_of_longer_match = False
+ else:
+ is_prefix_of_longer_match = self._is_prefix_of_longer_match(buffer)
+
+ # When eager matches were found, give priority to them and also
+ # ignore all the longer matches.
+ eager_matches = [m for m in matches if m.eager()]
+
+ if eager_matches:
+ matches = eager_matches
+ is_prefix_of_longer_match = False
+
+ # Exact matches found, call handler.
+ if not is_prefix_of_longer_match and matches:
+ self._call_handler(matches[-1], key_sequence=buffer[:])
+ del buffer[:] # Keep reference.
+
+ # No match found.
+ elif not is_prefix_of_longer_match and not matches:
+ retry = True
+ found = False
+
+ # Loop over the input, try longest match first and shift.
+ for i in range(len(buffer), 0, -1):
+ matches = self._get_matches(buffer[:i])
+ if matches:
+ self._call_handler(matches[-1], key_sequence=buffer[:i])
+ del buffer[:i]
+ found = True
+ break
+
+ if not found:
+ del buffer[:1]
+
+ def feed(self, key_press: KeyPress, first: bool = False) -> None:
+ """
+ Add a new :class:`KeyPress` to the input queue.
+ (Don't forget to call `process_keys` in order to process the queue.)
+
+ :param first: If true, insert before everything else.
+ """
+ if first:
+ self.input_queue.appendleft(key_press)
+ else:
+ self.input_queue.append(key_press)
+
+ def feed_multiple(self, key_presses: list[KeyPress], first: bool = False) -> None:
+ """
+ :param first: If true, insert before everything else.
+ """
+ if first:
+ self.input_queue.extendleft(reversed(key_presses))
+ else:
+ self.input_queue.extend(key_presses)
+
+ def process_keys(self) -> None:
+ """
+ Process all the keys in the `input_queue`.
+ (To be called after `feed`.)
+
+ Note: because of the `feed`/`process_keys` separation, it is
+ possible to call `feed` from inside a key binding.
+ This function keeps looping until the queue is empty.
+ """
+ app = get_app()
+
+ def not_empty() -> bool:
+ # When the application result is set, stop processing keys. (E.g.
+ # if ENTER was received, followed by a few additional key strokes,
+ # leave the other keys in the queue.)
+ if app.is_done:
+ # But if there are still CPRResponse keys in the queue, these
+ # need to be processed.
+ return any(k for k in self.input_queue if k.key == Keys.CPRResponse)
+ else:
+ return bool(self.input_queue)
+
+ def get_next() -> KeyPress:
+ if app.is_done:
+ # Only process CPR responses. Everything else is typeahead.
+ cpr = [k for k in self.input_queue if k.key == Keys.CPRResponse][0]
+ self.input_queue.remove(cpr)
+ return cpr
+ else:
+ return self.input_queue.popleft()
+
+ is_flush = False
+
+ while not_empty():
+ # Process next key.
+ key_press = get_next()
+
+ is_flush = key_press is _Flush
+ is_cpr = key_press.key == Keys.CPRResponse
+
+ if not is_flush and not is_cpr:
+ self.before_key_press.fire()
+
+ try:
+ self._process_coroutine.send(key_press)
+ except Exception:
+ # If for some reason something goes wrong in the parser, (maybe
+ # an exception was raised) restart the processor for next time.
+ self.reset()
+ self.empty_queue()
+ raise
+
+ if not is_flush and not is_cpr:
+ self.after_key_press.fire()
+
+ # Skip timeout if the last key was flush.
+ if not is_flush:
+ self._start_timeout()
+
+ def empty_queue(self) -> list[KeyPress]:
+ """
+ Empty the input queue. Return the unprocessed input.
+ """
+ key_presses = list(self.input_queue)
+ self.input_queue.clear()
+
+ # Filter out CPRs. We don't want to return these.
+ key_presses = [k for k in key_presses if k.key != Keys.CPRResponse]
+ return key_presses
+
+ def _call_handler(self, handler: Binding, key_sequence: list[KeyPress]) -> None:
+ app = get_app()
+ was_recording_emacs = app.emacs_state.is_recording
+ was_recording_vi = bool(app.vi_state.recording_register)
+ was_temporary_navigation_mode = app.vi_state.temporary_navigation_mode
+ arg = self.arg
+ self.arg = None
+
+ event = KeyPressEvent(
+ weakref.ref(self),
+ arg=arg,
+ key_sequence=key_sequence,
+ previous_key_sequence=self._previous_key_sequence,
+ is_repeat=(handler == self._previous_handler),
+ )
+
+ # Save the state of the current buffer.
+ if handler.save_before(event):
+ event.app.current_buffer.save_to_undo_stack()
+
+ # Call handler.
+ from prompt_toolkit.buffer import EditReadOnlyBuffer
+
+ try:
+ handler.call(event)
+ self._fix_vi_cursor_position(event)
+
+ except EditReadOnlyBuffer:
+ # When a key binding does an attempt to change a buffer which is
+ # read-only, we can ignore that. We sound a bell and go on.
+ app.output.bell()
+
+ if was_temporary_navigation_mode:
+ self._leave_vi_temp_navigation_mode(event)
+
+ self._previous_key_sequence = key_sequence
+ self._previous_handler = handler
+
+ # Record the key sequence in our macro. (Only if we're in macro mode
+ # before and after executing the key.)
+ if handler.record_in_macro():
+ if app.emacs_state.is_recording and was_recording_emacs:
+ recording = app.emacs_state.current_recording
+ if recording is not None: # Should always be true, given that
+ # `was_recording_emacs` is set.
+ recording.extend(key_sequence)
+
+ if app.vi_state.recording_register and was_recording_vi:
+ for k in key_sequence:
+ app.vi_state.current_recording += k.data
+
+ def _fix_vi_cursor_position(self, event: KeyPressEvent) -> None:
+ """
+ After every command, make sure that if we are in Vi navigation mode, we
+ never put the cursor after the last character of a line. (Unless it's
+ an empty line.)
+ """
+ app = event.app
+ buff = app.current_buffer
+ preferred_column = buff.preferred_column
+
+ if (
+ vi_navigation_mode()
+ and buff.document.is_cursor_at_the_end_of_line
+ and len(buff.document.current_line) > 0
+ ):
+ buff.cursor_position -= 1
+
+ # Set the preferred_column for arrow up/down again.
+ # (This was cleared after changing the cursor position.)
+ buff.preferred_column = preferred_column
+
+ def _leave_vi_temp_navigation_mode(self, event: KeyPressEvent) -> None:
+ """
+ If we're in Vi temporary navigation (normal) mode, return to
+ insert/replace mode after executing one action.
+ """
+ app = event.app
+
+ if app.editing_mode == EditingMode.VI:
+ # Not waiting for a text object and no argument has been given.
+ if app.vi_state.operator_func is None and self.arg is None:
+ app.vi_state.temporary_navigation_mode = False
+
+ def _start_timeout(self) -> None:
+ """
+ Start auto flush timeout. Similar to Vim's `timeoutlen` option.
+
+ Start a background coroutine with a timer. When this timeout expires
+ and no key was pressed in the meantime, we flush all data in the queue
+ and call the appropriate key binding handlers.
+ """
+ app = get_app()
+ timeout = app.timeoutlen
+
+ if timeout is None:
+ return
+
+ async def wait() -> None:
+ "Wait for timeout."
+ # This sleep can be cancelled. In that case we don't flush.
+ await sleep(timeout)
+
+ if len(self.key_buffer) > 0:
+ # (No keys pressed in the meantime.)
+ flush_keys()
+
+ def flush_keys() -> None:
+ "Flush keys."
+ self.feed(_Flush)
+ self.process_keys()
+
+ # Automatically flush keys.
+ if self._flush_wait_task:
+ self._flush_wait_task.cancel()
+ self._flush_wait_task = app.create_background_task(wait())
+
+ def send_sigint(self) -> None:
+ """
+ Send SIGINT. Immediately call the SIGINT key handler.
+ """
+ self.feed(KeyPress(key=Keys.SIGINT), first=True)
+ self.process_keys()
+
+
+class KeyPressEvent:
+ """
+ Key press event, delivered to key bindings.
+
+ :param key_processor_ref: Weak reference to the `KeyProcessor`.
+ :param arg: Repetition argument.
+ :param key_sequence: List of `KeyPress` instances.
+ :param previouskey_sequence: Previous list of `KeyPress` instances.
+ :param is_repeat: True when the previous event was delivered to the same handler.
+ """
+
+ def __init__(
+ self,
+ key_processor_ref: weakref.ReferenceType[KeyProcessor],
+ arg: str | None,
+ key_sequence: list[KeyPress],
+ previous_key_sequence: list[KeyPress],
+ is_repeat: bool,
+ ) -> None:
+ self._key_processor_ref = key_processor_ref
+ self.key_sequence = key_sequence
+ self.previous_key_sequence = previous_key_sequence
+
+ #: True when the previous key sequence was handled by the same handler.
+ self.is_repeat = is_repeat
+
+ self._arg = arg
+ self._app = get_app()
+
+ def __repr__(self) -> str:
+ return f"KeyPressEvent(arg={self.arg!r}, key_sequence={self.key_sequence!r}, is_repeat={self.is_repeat!r})"
+
+ @property
+ def data(self) -> str:
+ return self.key_sequence[-1].data
+
+ @property
+ def key_processor(self) -> KeyProcessor:
+ processor = self._key_processor_ref()
+ if processor is None:
+ raise Exception("KeyProcessor was lost. This should not happen.")
+ return processor
+
+ @property
+ def app(self) -> Application[Any]:
+ """
+ The current `Application` object.
+ """
+ return self._app
+
+ @property
+ def current_buffer(self) -> Buffer:
+ """
+ The current buffer.
+ """
+ return self.app.current_buffer
+
+ @property
+ def arg(self) -> int:
+ """
+ Repetition argument.
+ """
+ if self._arg == "-":
+ return -1
+
+ result = int(self._arg or 1)
+
+ # Don't exceed a million.
+ if int(result) >= 1000000:
+ result = 1
+
+ return result
+
+ @property
+ def arg_present(self) -> bool:
+ """
+ True if repetition argument was explicitly provided.
+ """
+ return self._arg is not None
+
+ def append_to_arg_count(self, data: str) -> None:
+ """
+ Add digit to the input argument.
+
+ :param data: the typed digit as string
+ """
+ assert data in "-0123456789"
+ current = self._arg
+
+ if data == "-":
+ assert current is None or current == "-"
+ result = data
+ elif current is None:
+ result = data
+ else:
+ result = f"{current}{data}"
+
+ self.key_processor.arg = result
+
+ @property
+ def cli(self) -> Application[Any]:
+ "For backward-compatibility."
+ return self.app
diff --git a/lib/python3.12/site-packages/prompt_toolkit/key_binding/vi_state.py b/lib/python3.12/site-packages/prompt_toolkit/key_binding/vi_state.py
new file mode 100644
index 0000000000000000000000000000000000000000..0543911e3285ee7f9e0fd763b8e0a56566bb538e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/key_binding/vi_state.py
@@ -0,0 +1,107 @@
+from __future__ import annotations
+
+from enum import Enum
+from typing import TYPE_CHECKING, Callable
+
+from prompt_toolkit.clipboard import ClipboardData
+
+if TYPE_CHECKING:
+ from .bindings.vi import TextObject
+ from .key_processor import KeyPressEvent
+
+__all__ = [
+ "InputMode",
+ "CharacterFind",
+ "ViState",
+]
+
+
+class InputMode(str, Enum):
+ value: str
+
+ INSERT = "vi-insert"
+ INSERT_MULTIPLE = "vi-insert-multiple"
+ NAVIGATION = "vi-navigation" # Normal mode.
+ REPLACE = "vi-replace"
+ REPLACE_SINGLE = "vi-replace-single"
+
+
+class CharacterFind:
+ def __init__(self, character: str, backwards: bool = False) -> None:
+ self.character = character
+ self.backwards = backwards
+
+
+class ViState:
+ """
+ Mutable class to hold the state of the Vi navigation.
+ """
+
+ def __init__(self) -> None:
+ #: None or CharacterFind instance. (This is used to repeat the last
+ #: search in Vi mode, by pressing the 'n' or 'N' in navigation mode.)
+ self.last_character_find: CharacterFind | None = None
+
+ # When an operator is given and we are waiting for text object,
+ # -- e.g. in the case of 'dw', after the 'd' --, an operator callback
+ # is set here.
+ self.operator_func: None | (Callable[[KeyPressEvent, TextObject], None]) = None
+ self.operator_arg: int | None = None
+
+ #: Named registers. Maps register name (e.g. 'a') to
+ #: :class:`ClipboardData` instances.
+ self.named_registers: dict[str, ClipboardData] = {}
+
+ #: The Vi mode we're currently in to.
+ self.__input_mode = InputMode.INSERT
+
+ #: Waiting for digraph.
+ self.waiting_for_digraph = False
+ self.digraph_symbol1: str | None = None # (None or a symbol.)
+
+ #: When true, make ~ act as an operator.
+ self.tilde_operator = False
+
+ #: Register in which we are recording a macro.
+ #: `None` when not recording anything.
+ # Note that the recording is only stored in the register after the
+ # recording is stopped. So we record in a separate `current_recording`
+ # variable.
+ self.recording_register: str | None = None
+ self.current_recording: str = ""
+
+ # Temporary navigation (normal) mode.
+ # This happens when control-o has been pressed in insert or replace
+ # mode. The user can now do one navigation action and we'll return back
+ # to insert/replace.
+ self.temporary_navigation_mode = False
+
+ @property
+ def input_mode(self) -> InputMode:
+ "Get `InputMode`."
+ return self.__input_mode
+
+ @input_mode.setter
+ def input_mode(self, value: InputMode) -> None:
+ "Set `InputMode`."
+ if value == InputMode.NAVIGATION:
+ self.waiting_for_digraph = False
+ self.operator_func = None
+ self.operator_arg = None
+
+ self.__input_mode = value
+
+ def reset(self) -> None:
+ """
+ Reset state, go back to the given mode. INSERT by default.
+ """
+ # Go back to insert mode.
+ self.input_mode = InputMode.INSERT
+
+ self.waiting_for_digraph = False
+ self.operator_func = None
+ self.operator_arg = None
+
+ # Reset recording state.
+ self.recording_register = None
+ self.current_recording = ""
diff --git a/lib/python3.12/site-packages/prompt_toolkit/keys.py b/lib/python3.12/site-packages/prompt_toolkit/keys.py
new file mode 100644
index 0000000000000000000000000000000000000000..ee52aee86a1ce1812d3326945edd223966ffedfe
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/keys.py
@@ -0,0 +1,222 @@
+from __future__ import annotations
+
+from enum import Enum
+
+__all__ = [
+ "Keys",
+ "ALL_KEYS",
+]
+
+
+class Keys(str, Enum):
+ """
+ List of keys for use in key bindings.
+
+ Note that this is an "StrEnum", all values can be compared against
+ strings.
+ """
+
+ value: str
+
+ Escape = "escape" # Also Control-[
+ ShiftEscape = "s-escape"
+
+ ControlAt = "c-@" # Also Control-Space.
+
+ ControlA = "c-a"
+ ControlB = "c-b"
+ ControlC = "c-c"
+ ControlD = "c-d"
+ ControlE = "c-e"
+ ControlF = "c-f"
+ ControlG = "c-g"
+ ControlH = "c-h"
+ ControlI = "c-i" # Tab
+ ControlJ = "c-j" # Newline
+ ControlK = "c-k"
+ ControlL = "c-l"
+ ControlM = "c-m" # Carriage return
+ ControlN = "c-n"
+ ControlO = "c-o"
+ ControlP = "c-p"
+ ControlQ = "c-q"
+ ControlR = "c-r"
+ ControlS = "c-s"
+ ControlT = "c-t"
+ ControlU = "c-u"
+ ControlV = "c-v"
+ ControlW = "c-w"
+ ControlX = "c-x"
+ ControlY = "c-y"
+ ControlZ = "c-z"
+
+ Control1 = "c-1"
+ Control2 = "c-2"
+ Control3 = "c-3"
+ Control4 = "c-4"
+ Control5 = "c-5"
+ Control6 = "c-6"
+ Control7 = "c-7"
+ Control8 = "c-8"
+ Control9 = "c-9"
+ Control0 = "c-0"
+
+ ControlShift1 = "c-s-1"
+ ControlShift2 = "c-s-2"
+ ControlShift3 = "c-s-3"
+ ControlShift4 = "c-s-4"
+ ControlShift5 = "c-s-5"
+ ControlShift6 = "c-s-6"
+ ControlShift7 = "c-s-7"
+ ControlShift8 = "c-s-8"
+ ControlShift9 = "c-s-9"
+ ControlShift0 = "c-s-0"
+
+ ControlBackslash = "c-\\"
+ ControlSquareClose = "c-]"
+ ControlCircumflex = "c-^"
+ ControlUnderscore = "c-_"
+
+ Left = "left"
+ Right = "right"
+ Up = "up"
+ Down = "down"
+ Home = "home"
+ End = "end"
+ Insert = "insert"
+ Delete = "delete"
+ PageUp = "pageup"
+ PageDown = "pagedown"
+
+ ControlLeft = "c-left"
+ ControlRight = "c-right"
+ ControlUp = "c-up"
+ ControlDown = "c-down"
+ ControlHome = "c-home"
+ ControlEnd = "c-end"
+ ControlInsert = "c-insert"
+ ControlDelete = "c-delete"
+ ControlPageUp = "c-pageup"
+ ControlPageDown = "c-pagedown"
+
+ ShiftLeft = "s-left"
+ ShiftRight = "s-right"
+ ShiftUp = "s-up"
+ ShiftDown = "s-down"
+ ShiftHome = "s-home"
+ ShiftEnd = "s-end"
+ ShiftInsert = "s-insert"
+ ShiftDelete = "s-delete"
+ ShiftPageUp = "s-pageup"
+ ShiftPageDown = "s-pagedown"
+
+ ControlShiftLeft = "c-s-left"
+ ControlShiftRight = "c-s-right"
+ ControlShiftUp = "c-s-up"
+ ControlShiftDown = "c-s-down"
+ ControlShiftHome = "c-s-home"
+ ControlShiftEnd = "c-s-end"
+ ControlShiftInsert = "c-s-insert"
+ ControlShiftDelete = "c-s-delete"
+ ControlShiftPageUp = "c-s-pageup"
+ ControlShiftPageDown = "c-s-pagedown"
+
+ BackTab = "s-tab" # shift + tab
+
+ F1 = "f1"
+ F2 = "f2"
+ F3 = "f3"
+ F4 = "f4"
+ F5 = "f5"
+ F6 = "f6"
+ F7 = "f7"
+ F8 = "f8"
+ F9 = "f9"
+ F10 = "f10"
+ F11 = "f11"
+ F12 = "f12"
+ F13 = "f13"
+ F14 = "f14"
+ F15 = "f15"
+ F16 = "f16"
+ F17 = "f17"
+ F18 = "f18"
+ F19 = "f19"
+ F20 = "f20"
+ F21 = "f21"
+ F22 = "f22"
+ F23 = "f23"
+ F24 = "f24"
+
+ ControlF1 = "c-f1"
+ ControlF2 = "c-f2"
+ ControlF3 = "c-f3"
+ ControlF4 = "c-f4"
+ ControlF5 = "c-f5"
+ ControlF6 = "c-f6"
+ ControlF7 = "c-f7"
+ ControlF8 = "c-f8"
+ ControlF9 = "c-f9"
+ ControlF10 = "c-f10"
+ ControlF11 = "c-f11"
+ ControlF12 = "c-f12"
+ ControlF13 = "c-f13"
+ ControlF14 = "c-f14"
+ ControlF15 = "c-f15"
+ ControlF16 = "c-f16"
+ ControlF17 = "c-f17"
+ ControlF18 = "c-f18"
+ ControlF19 = "c-f19"
+ ControlF20 = "c-f20"
+ ControlF21 = "c-f21"
+ ControlF22 = "c-f22"
+ ControlF23 = "c-f23"
+ ControlF24 = "c-f24"
+
+ # Matches any key.
+ Any = ""
+
+ # Special.
+ ScrollUp = ""
+ ScrollDown = ""
+
+ CPRResponse = ""
+ Vt100MouseEvent = ""
+ WindowsMouseEvent = ""
+ BracketedPaste = ""
+
+ SIGINT = ""
+
+ # For internal use: key which is ignored.
+ # (The key binding for this key should not do anything.)
+ Ignore = ""
+
+ # Some 'Key' aliases (for backwards-compatibility).
+ ControlSpace = ControlAt
+ Tab = ControlI
+ Enter = ControlM
+ Backspace = ControlH
+
+ # ShiftControl was renamed to ControlShift in
+ # 888fcb6fa4efea0de8333177e1bbc792f3ff3c24 (20 Feb 2020).
+ ShiftControlLeft = ControlShiftLeft
+ ShiftControlRight = ControlShiftRight
+ ShiftControlHome = ControlShiftHome
+ ShiftControlEnd = ControlShiftEnd
+
+
+ALL_KEYS: list[str] = [k.value for k in Keys]
+
+
+# Aliases.
+KEY_ALIASES: dict[str, str] = {
+ "backspace": "c-h",
+ "c-space": "c-@",
+ "enter": "c-m",
+ "tab": "c-i",
+ # ShiftControl was renamed to ControlShift.
+ "s-c-left": "c-s-left",
+ "s-c-right": "c-s-right",
+ "s-c-home": "c-s-home",
+ "s-c-end": "c-s-end",
+}
diff --git a/lib/python3.12/site-packages/prompt_toolkit/log.py b/lib/python3.12/site-packages/prompt_toolkit/log.py
new file mode 100644
index 0000000000000000000000000000000000000000..285357927ad9eb15e753bbe90445e798e94f7b99
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/log.py
@@ -0,0 +1,13 @@
+"""
+Logging configuration.
+"""
+
+from __future__ import annotations
+
+import logging
+
+__all__ = [
+ "logger",
+]
+
+logger = logging.getLogger(__package__)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/mouse_events.py b/lib/python3.12/site-packages/prompt_toolkit/mouse_events.py
new file mode 100644
index 0000000000000000000000000000000000000000..f244f8eac42f687eaf67ba720d5b7fda9970477a
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/mouse_events.py
@@ -0,0 +1,85 @@
+"""
+Mouse events.
+
+
+How it works
+------------
+
+The renderer has a 2 dimensional grid of mouse event handlers.
+(`prompt_toolkit.layout.MouseHandlers`.) When the layout is rendered, the
+`Window` class will make sure that this grid will also be filled with
+callbacks. For vt100 terminals, mouse events are received through stdin, just
+like any other key press. There is a handler among the key bindings that
+catches these events and forwards them to such a mouse event handler. It passes
+through the `Window` class where the coordinates are translated from absolute
+coordinates to coordinates relative to the user control, and there
+`UIControl.mouse_handler` is called.
+"""
+
+from __future__ import annotations
+
+from enum import Enum
+
+from .data_structures import Point
+
+__all__ = ["MouseEventType", "MouseButton", "MouseModifier", "MouseEvent"]
+
+
+class MouseEventType(Enum):
+ # Mouse up: This same event type is fired for all three events: left mouse
+ # up, right mouse up, or middle mouse up
+ MOUSE_UP = "MOUSE_UP"
+
+ # Mouse down: This implicitly refers to the left mouse down (this event is
+ # not fired upon pressing the middle or right mouse buttons).
+ MOUSE_DOWN = "MOUSE_DOWN"
+
+ SCROLL_UP = "SCROLL_UP"
+ SCROLL_DOWN = "SCROLL_DOWN"
+
+ # Triggered when the left mouse button is held down, and the mouse moves
+ MOUSE_MOVE = "MOUSE_MOVE"
+
+
+class MouseButton(Enum):
+ LEFT = "LEFT"
+ MIDDLE = "MIDDLE"
+ RIGHT = "RIGHT"
+
+ # When we're scrolling, or just moving the mouse and not pressing a button.
+ NONE = "NONE"
+
+ # This is for when we don't know which mouse button was pressed, but we do
+ # know that one has been pressed during this mouse event (as opposed to
+ # scrolling, for example)
+ UNKNOWN = "UNKNOWN"
+
+
+class MouseModifier(Enum):
+ SHIFT = "SHIFT"
+ ALT = "ALT"
+ CONTROL = "CONTROL"
+
+
+class MouseEvent:
+ """
+ Mouse event, sent to `UIControl.mouse_handler`.
+
+ :param position: `Point` instance.
+ :param event_type: `MouseEventType`.
+ """
+
+ def __init__(
+ self,
+ position: Point,
+ event_type: MouseEventType,
+ button: MouseButton,
+ modifiers: frozenset[MouseModifier],
+ ) -> None:
+ self.position = position
+ self.event_type = event_type
+ self.button = button
+ self.modifiers = modifiers
+
+ def __repr__(self) -> str:
+ return f"MouseEvent({self.position!r},{self.event_type!r},{self.button!r},{self.modifiers!r})"
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/output/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..6b4c5f3b66c0032074d35edf3edec896f340888d
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/__init__.py
@@ -0,0 +1,15 @@
+from __future__ import annotations
+
+from .base import DummyOutput, Output
+from .color_depth import ColorDepth
+from .defaults import create_output
+
+__all__ = [
+ # Base.
+ "Output",
+ "DummyOutput",
+ # Color depth.
+ "ColorDepth",
+ # Defaults.
+ "create_output",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..e2ef766dca05193bbd7fdd96435f02638461d3a5
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/base.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/base.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..a4194845ca00b61ffb6050a28a8b3b8a9a3509c0
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/base.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/color_depth.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/color_depth.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..abba77c120945e4b3cfc58e88eb40f079a7d1ba2
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/color_depth.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/conemu.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/conemu.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..380da634cb9aaad4a330174bc07392b4f7e116c0
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/conemu.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/defaults.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/defaults.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..1514c5f7fee974f4d41f822cf3cf1c1eaad61f64
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/defaults.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/flush_stdout.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/flush_stdout.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..d36569673af0e11394b25676fad217655799ed93
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/flush_stdout.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/plain_text.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/plain_text.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..9875767f4d10a0b479b673d9e60e8e3d561edc75
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/plain_text.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/vt100.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/vt100.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..b10bc04b0303ca947c688ce9050a9c7983b5278f
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/vt100.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/win32.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/win32.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..f4c29f83d9da7405ca0bcaa8dd4dc984bf5211ce
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/win32.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/windows10.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/windows10.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..9dc7593abde46a9694ad20c804507047d4417f61
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/output/__pycache__/windows10.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/base.py b/lib/python3.12/site-packages/prompt_toolkit/output/base.py
new file mode 100644
index 0000000000000000000000000000000000000000..6ba09fdd089f02e047d95e2ea77a9d7d3e835728
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/base.py
@@ -0,0 +1,332 @@
+"""
+Interface for an output.
+"""
+
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from typing import TextIO
+
+from prompt_toolkit.cursor_shapes import CursorShape
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.styles import Attrs
+
+from .color_depth import ColorDepth
+
+__all__ = [
+ "Output",
+ "DummyOutput",
+]
+
+
+class Output(metaclass=ABCMeta):
+ """
+ Base class defining the output interface for a
+ :class:`~prompt_toolkit.renderer.Renderer`.
+
+ Actual implementations are
+ :class:`~prompt_toolkit.output.vt100.Vt100_Output` and
+ :class:`~prompt_toolkit.output.win32.Win32Output`.
+ """
+
+ stdout: TextIO | None = None
+
+ @abstractmethod
+ def fileno(self) -> int:
+ "Return the file descriptor to which we can write for the output."
+
+ @abstractmethod
+ def encoding(self) -> str:
+ """
+ Return the encoding for this output, e.g. 'utf-8'.
+ (This is used mainly to know which characters are supported by the
+ output the data, so that the UI can provide alternatives, when
+ required.)
+ """
+
+ @abstractmethod
+ def write(self, data: str) -> None:
+ "Write text (Terminal escape sequences will be removed/escaped.)"
+
+ @abstractmethod
+ def write_raw(self, data: str) -> None:
+ "Write text."
+
+ @abstractmethod
+ def set_title(self, title: str) -> None:
+ "Set terminal title."
+
+ @abstractmethod
+ def clear_title(self) -> None:
+ "Clear title again. (or restore previous title.)"
+
+ @abstractmethod
+ def flush(self) -> None:
+ "Write to output stream and flush."
+
+ @abstractmethod
+ def erase_screen(self) -> None:
+ """
+ Erases the screen with the background color and moves the cursor to
+ home.
+ """
+
+ @abstractmethod
+ def enter_alternate_screen(self) -> None:
+ "Go to the alternate screen buffer. (For full screen applications)."
+
+ @abstractmethod
+ def quit_alternate_screen(self) -> None:
+ "Leave the alternate screen buffer."
+
+ @abstractmethod
+ def enable_mouse_support(self) -> None:
+ "Enable mouse."
+
+ @abstractmethod
+ def disable_mouse_support(self) -> None:
+ "Disable mouse."
+
+ @abstractmethod
+ def erase_end_of_line(self) -> None:
+ """
+ Erases from the current cursor position to the end of the current line.
+ """
+
+ @abstractmethod
+ def erase_down(self) -> None:
+ """
+ Erases the screen from the current line down to the bottom of the
+ screen.
+ """
+
+ @abstractmethod
+ def reset_attributes(self) -> None:
+ "Reset color and styling attributes."
+
+ @abstractmethod
+ def set_attributes(self, attrs: Attrs, color_depth: ColorDepth) -> None:
+ "Set new color and styling attributes."
+
+ @abstractmethod
+ def disable_autowrap(self) -> None:
+ "Disable auto line wrapping."
+
+ @abstractmethod
+ def enable_autowrap(self) -> None:
+ "Enable auto line wrapping."
+
+ @abstractmethod
+ def cursor_goto(self, row: int = 0, column: int = 0) -> None:
+ "Move cursor position."
+
+ @abstractmethod
+ def cursor_up(self, amount: int) -> None:
+ "Move cursor `amount` place up."
+
+ @abstractmethod
+ def cursor_down(self, amount: int) -> None:
+ "Move cursor `amount` place down."
+
+ @abstractmethod
+ def cursor_forward(self, amount: int) -> None:
+ "Move cursor `amount` place forward."
+
+ @abstractmethod
+ def cursor_backward(self, amount: int) -> None:
+ "Move cursor `amount` place backward."
+
+ @abstractmethod
+ def hide_cursor(self) -> None:
+ "Hide cursor."
+
+ @abstractmethod
+ def show_cursor(self) -> None:
+ "Show cursor."
+
+ @abstractmethod
+ def set_cursor_shape(self, cursor_shape: CursorShape) -> None:
+ "Set cursor shape to block, beam or underline."
+
+ @abstractmethod
+ def reset_cursor_shape(self) -> None:
+ "Reset cursor shape."
+
+ def ask_for_cpr(self) -> None:
+ """
+ Asks for a cursor position report (CPR).
+ (VT100 only.)
+ """
+
+ @property
+ def responds_to_cpr(self) -> bool:
+ """
+ `True` if the `Application` can expect to receive a CPR response after
+ calling `ask_for_cpr` (this will come back through the corresponding
+ `Input`).
+
+ This is used to determine the amount of available rows we have below
+ the cursor position. In the first place, we have this so that the drop
+ down autocompletion menus are sized according to the available space.
+
+ On Windows, we don't need this, there we have
+ `get_rows_below_cursor_position`.
+ """
+ return False
+
+ @abstractmethod
+ def get_size(self) -> Size:
+ "Return the size of the output window."
+
+ def bell(self) -> None:
+ "Sound bell."
+
+ def enable_bracketed_paste(self) -> None:
+ "For vt100 only."
+
+ def disable_bracketed_paste(self) -> None:
+ "For vt100 only."
+
+ def reset_cursor_key_mode(self) -> None:
+ """
+ For vt100 only.
+ Put the terminal in normal cursor mode (instead of application mode).
+
+ See: https://vt100.net/docs/vt100-ug/chapter3.html
+ """
+
+ def scroll_buffer_to_prompt(self) -> None:
+ "For Win32 only."
+
+ def get_rows_below_cursor_position(self) -> int:
+ "For Windows only."
+ raise NotImplementedError
+
+ @abstractmethod
+ def get_default_color_depth(self) -> ColorDepth:
+ """
+ Get default color depth for this output.
+
+ This value will be used if no color depth was explicitly passed to the
+ `Application`.
+
+ .. note::
+
+ If the `$PROMPT_TOOLKIT_COLOR_DEPTH` environment variable has been
+ set, then `outputs.defaults.create_output` will pass this value to
+ the implementation as the default_color_depth, which is returned
+ here. (This is not used when the output corresponds to a
+ prompt_toolkit SSH/Telnet session.)
+ """
+
+
+class DummyOutput(Output):
+ """
+ For testing. An output class that doesn't render anything.
+ """
+
+ def fileno(self) -> int:
+ "There is no sensible default for fileno()."
+ raise NotImplementedError
+
+ def encoding(self) -> str:
+ return "utf-8"
+
+ def write(self, data: str) -> None:
+ pass
+
+ def write_raw(self, data: str) -> None:
+ pass
+
+ def set_title(self, title: str) -> None:
+ pass
+
+ def clear_title(self) -> None:
+ pass
+
+ def flush(self) -> None:
+ pass
+
+ def erase_screen(self) -> None:
+ pass
+
+ def enter_alternate_screen(self) -> None:
+ pass
+
+ def quit_alternate_screen(self) -> None:
+ pass
+
+ def enable_mouse_support(self) -> None:
+ pass
+
+ def disable_mouse_support(self) -> None:
+ pass
+
+ def erase_end_of_line(self) -> None:
+ pass
+
+ def erase_down(self) -> None:
+ pass
+
+ def reset_attributes(self) -> None:
+ pass
+
+ def set_attributes(self, attrs: Attrs, color_depth: ColorDepth) -> None:
+ pass
+
+ def disable_autowrap(self) -> None:
+ pass
+
+ def enable_autowrap(self) -> None:
+ pass
+
+ def cursor_goto(self, row: int = 0, column: int = 0) -> None:
+ pass
+
+ def cursor_up(self, amount: int) -> None:
+ pass
+
+ def cursor_down(self, amount: int) -> None:
+ pass
+
+ def cursor_forward(self, amount: int) -> None:
+ pass
+
+ def cursor_backward(self, amount: int) -> None:
+ pass
+
+ def hide_cursor(self) -> None:
+ pass
+
+ def show_cursor(self) -> None:
+ pass
+
+ def set_cursor_shape(self, cursor_shape: CursorShape) -> None:
+ pass
+
+ def reset_cursor_shape(self) -> None:
+ pass
+
+ def ask_for_cpr(self) -> None:
+ pass
+
+ def bell(self) -> None:
+ pass
+
+ def enable_bracketed_paste(self) -> None:
+ pass
+
+ def disable_bracketed_paste(self) -> None:
+ pass
+
+ def scroll_buffer_to_prompt(self) -> None:
+ pass
+
+ def get_size(self) -> Size:
+ return Size(rows=40, columns=80)
+
+ def get_rows_below_cursor_position(self) -> int:
+ return 40
+
+ def get_default_color_depth(self) -> ColorDepth:
+ return ColorDepth.DEPTH_1_BIT
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/color_depth.py b/lib/python3.12/site-packages/prompt_toolkit/output/color_depth.py
new file mode 100644
index 0000000000000000000000000000000000000000..f66d2bea301da9963918c28ba5ee277b755ceb5e
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/color_depth.py
@@ -0,0 +1,64 @@
+from __future__ import annotations
+
+import os
+from enum import Enum
+
+__all__ = [
+ "ColorDepth",
+]
+
+
+class ColorDepth(str, Enum):
+ """
+ Possible color depth values for the output.
+ """
+
+ value: str
+
+ #: One color only.
+ DEPTH_1_BIT = "DEPTH_1_BIT"
+
+ #: ANSI Colors.
+ DEPTH_4_BIT = "DEPTH_4_BIT"
+
+ #: The default.
+ DEPTH_8_BIT = "DEPTH_8_BIT"
+
+ #: 24 bit True color.
+ DEPTH_24_BIT = "DEPTH_24_BIT"
+
+ # Aliases.
+ MONOCHROME = DEPTH_1_BIT
+ ANSI_COLORS_ONLY = DEPTH_4_BIT
+ DEFAULT = DEPTH_8_BIT
+ TRUE_COLOR = DEPTH_24_BIT
+
+ @classmethod
+ def from_env(cls) -> ColorDepth | None:
+ """
+ Return the color depth if the $PROMPT_TOOLKIT_COLOR_DEPTH environment
+ variable has been set.
+
+ This is a way to enforce a certain color depth in all prompt_toolkit
+ applications.
+ """
+ # Disable color if a `NO_COLOR` environment variable is set.
+ # See: https://no-color.org/
+ if os.environ.get("NO_COLOR"):
+ return cls.DEPTH_1_BIT
+
+ # Check the `PROMPT_TOOLKIT_COLOR_DEPTH` environment variable.
+ all_values = [i.value for i in ColorDepth]
+ if os.environ.get("PROMPT_TOOLKIT_COLOR_DEPTH") in all_values:
+ return cls(os.environ["PROMPT_TOOLKIT_COLOR_DEPTH"])
+
+ return None
+
+ @classmethod
+ def default(cls) -> ColorDepth:
+ """
+ Return the default color depth for the default output.
+ """
+ from .defaults import create_output
+
+ return create_output().get_default_color_depth()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/conemu.py b/lib/python3.12/site-packages/prompt_toolkit/output/conemu.py
new file mode 100644
index 0000000000000000000000000000000000000000..636994494544ba223647bd1abbf33baaafedd89f
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/conemu.py
@@ -0,0 +1,65 @@
+from __future__ import annotations
+
+import sys
+
+assert sys.platform == "win32"
+
+from typing import Any, TextIO
+
+from prompt_toolkit.data_structures import Size
+
+from .base import Output
+from .color_depth import ColorDepth
+from .vt100 import Vt100_Output
+from .win32 import Win32Output
+
+__all__ = [
+ "ConEmuOutput",
+]
+
+
+class ConEmuOutput:
+ """
+ ConEmu (Windows) output abstraction.
+
+ ConEmu is a Windows console application, but it also supports ANSI escape
+ sequences. This output class is actually a proxy to both `Win32Output` and
+ `Vt100_Output`. It uses `Win32Output` for console sizing and scrolling, but
+ all cursor movements and scrolling happens through the `Vt100_Output`.
+
+ This way, we can have 256 colors in ConEmu and Cmder. Rendering will be
+ even a little faster as well.
+
+ http://conemu.github.io/
+ http://gooseberrycreative.com/cmder/
+ """
+
+ def __init__(
+ self, stdout: TextIO, default_color_depth: ColorDepth | None = None
+ ) -> None:
+ self.win32_output = Win32Output(stdout, default_color_depth=default_color_depth)
+ self.vt100_output = Vt100_Output(
+ stdout, lambda: Size(0, 0), default_color_depth=default_color_depth
+ )
+
+ @property
+ def responds_to_cpr(self) -> bool:
+ return False # We don't need this on Windows.
+
+ def __getattr__(self, name: str) -> Any:
+ if name in (
+ "get_size",
+ "get_rows_below_cursor_position",
+ "enable_mouse_support",
+ "disable_mouse_support",
+ "scroll_buffer_to_prompt",
+ "get_win32_screen_buffer_info",
+ "enable_bracketed_paste",
+ "disable_bracketed_paste",
+ ):
+ return getattr(self.win32_output, name)
+ else:
+ return getattr(self.vt100_output, name)
+
+
+Output.register(ConEmuOutput)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/defaults.py b/lib/python3.12/site-packages/prompt_toolkit/output/defaults.py
new file mode 100644
index 0000000000000000000000000000000000000000..6b06ed43c884e017bc01d53d8a4bd88617d2a741
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/defaults.py
@@ -0,0 +1,106 @@
+from __future__ import annotations
+
+import sys
+from typing import TYPE_CHECKING, TextIO, cast
+
+from prompt_toolkit.utils import (
+ get_bell_environment_variable,
+ get_term_environment_variable,
+ is_conemu_ansi,
+)
+
+from .base import DummyOutput, Output
+from .color_depth import ColorDepth
+from .plain_text import PlainTextOutput
+
+if TYPE_CHECKING:
+ from prompt_toolkit.patch_stdout import StdoutProxy
+
+
+__all__ = [
+ "create_output",
+]
+
+
+def create_output(
+ stdout: TextIO | StdoutProxy | None = None, always_prefer_tty: bool = False
+) -> Output:
+ """
+ Return an :class:`~prompt_toolkit.output.Output` instance for the command
+ line.
+
+ :param stdout: The stdout object
+ :param always_prefer_tty: When set, look for `sys.stderr` if `sys.stdout`
+ is not a TTY. Useful if `sys.stdout` is redirected to a file, but we
+ still want user input and output on the terminal.
+
+ By default, this is `False`. If `sys.stdout` is not a terminal (maybe
+ it's redirected to a file), then a `PlainTextOutput` will be returned.
+ That way, tools like `print_formatted_text` will write plain text into
+ that file.
+ """
+ # Consider TERM, PROMPT_TOOLKIT_BELL, and PROMPT_TOOLKIT_COLOR_DEPTH
+ # environment variables. Notice that PROMPT_TOOLKIT_COLOR_DEPTH value is
+ # the default that's used if the Application doesn't override it.
+ term_from_env = get_term_environment_variable()
+ bell_from_env = get_bell_environment_variable()
+ color_depth_from_env = ColorDepth.from_env()
+
+ if stdout is None:
+ # By default, render to stdout. If the output is piped somewhere else,
+ # render to stderr.
+ stdout = sys.stdout
+
+ if always_prefer_tty:
+ for io in [sys.stdout, sys.stderr]:
+ if io is not None and io.isatty():
+ # (This is `None` when using `pythonw.exe` on Windows.)
+ stdout = io
+ break
+
+ # If the patch_stdout context manager has been used, then sys.stdout is
+ # replaced by this proxy. For prompt_toolkit applications, we want to use
+ # the real stdout.
+ from prompt_toolkit.patch_stdout import StdoutProxy
+
+ while isinstance(stdout, StdoutProxy):
+ stdout = stdout.original_stdout
+
+ # If the output is still `None`, use a DummyOutput.
+ # This happens for instance on Windows, when running the application under
+ # `pythonw.exe`. In that case, there won't be a terminal Window, and
+ # stdin/stdout/stderr are `None`.
+ if stdout is None:
+ return DummyOutput()
+
+ if sys.platform == "win32":
+ from .conemu import ConEmuOutput
+ from .win32 import Win32Output
+ from .windows10 import Windows10_Output, is_win_vt100_enabled
+
+ if is_win_vt100_enabled():
+ return cast(
+ Output,
+ Windows10_Output(stdout, default_color_depth=color_depth_from_env),
+ )
+ if is_conemu_ansi():
+ return cast(
+ Output, ConEmuOutput(stdout, default_color_depth=color_depth_from_env)
+ )
+ else:
+ return Win32Output(stdout, default_color_depth=color_depth_from_env)
+ else:
+ from .vt100 import Vt100_Output
+
+ # Stdout is not a TTY? Render as plain text.
+ # This is mostly useful if stdout is redirected to a file, and
+ # `print_formatted_text` is used.
+ if not stdout.isatty():
+ return PlainTextOutput(stdout)
+
+ return Vt100_Output.from_pty(
+ stdout,
+ term=term_from_env,
+ default_color_depth=color_depth_from_env,
+ enable_bell=bell_from_env,
+ )
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/flush_stdout.py b/lib/python3.12/site-packages/prompt_toolkit/output/flush_stdout.py
new file mode 100644
index 0000000000000000000000000000000000000000..daf58efee63cc6cbe10e9b522a2dc50394acd83f
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/flush_stdout.py
@@ -0,0 +1,87 @@
+from __future__ import annotations
+
+import errno
+import os
+import sys
+from contextlib import contextmanager
+from typing import IO, Iterator, TextIO
+
+__all__ = ["flush_stdout"]
+
+
+def flush_stdout(stdout: TextIO, data: str) -> None:
+ # If the IO object has an `encoding` and `buffer` attribute, it means that
+ # we can access the underlying BinaryIO object and write into it in binary
+ # mode. This is preferred if possible.
+ # NOTE: When used in a Jupyter notebook, don't write binary.
+ # `ipykernel.iostream.OutStream` has an `encoding` attribute, but not
+ # a `buffer` attribute, so we can't write binary in it.
+ has_binary_io = hasattr(stdout, "encoding") and hasattr(stdout, "buffer")
+
+ try:
+ # Ensure that `stdout` is made blocking when writing into it.
+ # Otherwise, when uvloop is activated (which makes stdout
+ # non-blocking), and we write big amounts of text, then we get a
+ # `BlockingIOError` here.
+ with _blocking_io(stdout):
+ # (We try to encode ourself, because that way we can replace
+ # characters that don't exist in the character set, avoiding
+ # UnicodeEncodeError crashes. E.g. u'\xb7' does not appear in 'ascii'.)
+ # My Arch Linux installation of july 2015 reported 'ANSI_X3.4-1968'
+ # for sys.stdout.encoding in xterm.
+ if has_binary_io:
+ stdout.buffer.write(data.encode(stdout.encoding or "utf-8", "replace"))
+ else:
+ stdout.write(data)
+
+ stdout.flush()
+ except OSError as e:
+ if e.args and e.args[0] == errno.EINTR:
+ # Interrupted system call. Can happen in case of a window
+ # resize signal. (Just ignore. The resize handler will render
+ # again anyway.)
+ pass
+ elif e.args and e.args[0] == 0:
+ # This can happen when there is a lot of output and the user
+ # sends a KeyboardInterrupt by pressing Control-C. E.g. in
+ # a Python REPL when we execute "while True: print('test')".
+ # (The `ptpython` REPL uses this `Output` class instead of
+ # `stdout` directly -- in order to be network transparent.)
+ # So, just ignore.
+ pass
+ else:
+ raise
+
+
+@contextmanager
+def _blocking_io(io: IO[str]) -> Iterator[None]:
+ """
+ Ensure that the FD for `io` is set to blocking in here.
+ """
+ if sys.platform == "win32":
+ # On Windows, the `os` module doesn't have a `get/set_blocking`
+ # function.
+ yield
+ return
+
+ try:
+ fd = io.fileno()
+ blocking = os.get_blocking(fd)
+ except: # noqa
+ # Failed somewhere.
+ # `get_blocking` can raise `OSError`.
+ # The io object can raise `AttributeError` when no `fileno()` method is
+ # present if we're not a real file object.
+ blocking = True # Assume we're good, and don't do anything.
+
+ try:
+ # Make blocking if we weren't blocking yet.
+ if not blocking:
+ os.set_blocking(fd, True)
+
+ yield
+
+ finally:
+ # Restore original blocking mode.
+ if not blocking:
+ os.set_blocking(fd, blocking)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/plain_text.py b/lib/python3.12/site-packages/prompt_toolkit/output/plain_text.py
new file mode 100644
index 0000000000000000000000000000000000000000..4b24ad960e4f2359e0363a216ad91a00d2910d24
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/plain_text.py
@@ -0,0 +1,143 @@
+from __future__ import annotations
+
+from typing import TextIO
+
+from prompt_toolkit.cursor_shapes import CursorShape
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.styles import Attrs
+
+from .base import Output
+from .color_depth import ColorDepth
+from .flush_stdout import flush_stdout
+
+__all__ = ["PlainTextOutput"]
+
+
+class PlainTextOutput(Output):
+ """
+ Output that won't include any ANSI escape sequences.
+
+ Useful when stdout is not a terminal. Maybe stdout is redirected to a file.
+ In this case, if `print_formatted_text` is used, for instance, we don't
+ want to include formatting.
+
+ (The code is mostly identical to `Vt100_Output`, but without the
+ formatting.)
+ """
+
+ def __init__(self, stdout: TextIO) -> None:
+ assert all(hasattr(stdout, a) for a in ("write", "flush"))
+
+ self.stdout: TextIO = stdout
+ self._buffer: list[str] = []
+
+ def fileno(self) -> int:
+ "There is no sensible default for fileno()."
+ return self.stdout.fileno()
+
+ def encoding(self) -> str:
+ return "utf-8"
+
+ def write(self, data: str) -> None:
+ self._buffer.append(data)
+
+ def write_raw(self, data: str) -> None:
+ self._buffer.append(data)
+
+ def set_title(self, title: str) -> None:
+ pass
+
+ def clear_title(self) -> None:
+ pass
+
+ def flush(self) -> None:
+ if not self._buffer:
+ return
+
+ data = "".join(self._buffer)
+ self._buffer = []
+ flush_stdout(self.stdout, data)
+
+ def erase_screen(self) -> None:
+ pass
+
+ def enter_alternate_screen(self) -> None:
+ pass
+
+ def quit_alternate_screen(self) -> None:
+ pass
+
+ def enable_mouse_support(self) -> None:
+ pass
+
+ def disable_mouse_support(self) -> None:
+ pass
+
+ def erase_end_of_line(self) -> None:
+ pass
+
+ def erase_down(self) -> None:
+ pass
+
+ def reset_attributes(self) -> None:
+ pass
+
+ def set_attributes(self, attrs: Attrs, color_depth: ColorDepth) -> None:
+ pass
+
+ def disable_autowrap(self) -> None:
+ pass
+
+ def enable_autowrap(self) -> None:
+ pass
+
+ def cursor_goto(self, row: int = 0, column: int = 0) -> None:
+ pass
+
+ def cursor_up(self, amount: int) -> None:
+ pass
+
+ def cursor_down(self, amount: int) -> None:
+ self._buffer.append("\n")
+
+ def cursor_forward(self, amount: int) -> None:
+ self._buffer.append(" " * amount)
+
+ def cursor_backward(self, amount: int) -> None:
+ pass
+
+ def hide_cursor(self) -> None:
+ pass
+
+ def show_cursor(self) -> None:
+ pass
+
+ def set_cursor_shape(self, cursor_shape: CursorShape) -> None:
+ pass
+
+ def reset_cursor_shape(self) -> None:
+ pass
+
+ def ask_for_cpr(self) -> None:
+ pass
+
+ def bell(self) -> None:
+ pass
+
+ def enable_bracketed_paste(self) -> None:
+ pass
+
+ def disable_bracketed_paste(self) -> None:
+ pass
+
+ def scroll_buffer_to_prompt(self) -> None:
+ pass
+
+ def get_size(self) -> Size:
+ return Size(rows=40, columns=80)
+
+ def get_rows_below_cursor_position(self) -> int:
+ return 8
+
+ def get_default_color_depth(self) -> ColorDepth:
+ return ColorDepth.DEPTH_1_BIT
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/vt100.py b/lib/python3.12/site-packages/prompt_toolkit/output/vt100.py
new file mode 100644
index 0000000000000000000000000000000000000000..57826b9409fcf3691a6191b803ef519ec0351ed9
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/vt100.py
@@ -0,0 +1,760 @@
+"""
+Output for vt100 terminals.
+
+A lot of thanks, regarding outputting of colors, goes to the Pygments project:
+(We don't rely on Pygments anymore, because many things are very custom, and
+everything has been highly optimized.)
+http://pygments.org/
+"""
+
+from __future__ import annotations
+
+import io
+import os
+import sys
+from typing import Callable, Dict, Hashable, Iterable, Sequence, TextIO, Tuple
+
+from prompt_toolkit.cursor_shapes import CursorShape
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.output import Output
+from prompt_toolkit.styles import ANSI_COLOR_NAMES, Attrs
+from prompt_toolkit.utils import is_dumb_terminal
+
+from .color_depth import ColorDepth
+from .flush_stdout import flush_stdout
+
+__all__ = [
+ "Vt100_Output",
+]
+
+
+FG_ANSI_COLORS = {
+ "ansidefault": 39,
+ # Low intensity.
+ "ansiblack": 30,
+ "ansired": 31,
+ "ansigreen": 32,
+ "ansiyellow": 33,
+ "ansiblue": 34,
+ "ansimagenta": 35,
+ "ansicyan": 36,
+ "ansigray": 37,
+ # High intensity.
+ "ansibrightblack": 90,
+ "ansibrightred": 91,
+ "ansibrightgreen": 92,
+ "ansibrightyellow": 93,
+ "ansibrightblue": 94,
+ "ansibrightmagenta": 95,
+ "ansibrightcyan": 96,
+ "ansiwhite": 97,
+}
+
+BG_ANSI_COLORS = {
+ "ansidefault": 49,
+ # Low intensity.
+ "ansiblack": 40,
+ "ansired": 41,
+ "ansigreen": 42,
+ "ansiyellow": 43,
+ "ansiblue": 44,
+ "ansimagenta": 45,
+ "ansicyan": 46,
+ "ansigray": 47,
+ # High intensity.
+ "ansibrightblack": 100,
+ "ansibrightred": 101,
+ "ansibrightgreen": 102,
+ "ansibrightyellow": 103,
+ "ansibrightblue": 104,
+ "ansibrightmagenta": 105,
+ "ansibrightcyan": 106,
+ "ansiwhite": 107,
+}
+
+
+ANSI_COLORS_TO_RGB = {
+ "ansidefault": (
+ 0x00,
+ 0x00,
+ 0x00,
+ ), # Don't use, 'default' doesn't really have a value.
+ "ansiblack": (0x00, 0x00, 0x00),
+ "ansigray": (0xE5, 0xE5, 0xE5),
+ "ansibrightblack": (0x7F, 0x7F, 0x7F),
+ "ansiwhite": (0xFF, 0xFF, 0xFF),
+ # Low intensity.
+ "ansired": (0xCD, 0x00, 0x00),
+ "ansigreen": (0x00, 0xCD, 0x00),
+ "ansiyellow": (0xCD, 0xCD, 0x00),
+ "ansiblue": (0x00, 0x00, 0xCD),
+ "ansimagenta": (0xCD, 0x00, 0xCD),
+ "ansicyan": (0x00, 0xCD, 0xCD),
+ # High intensity.
+ "ansibrightred": (0xFF, 0x00, 0x00),
+ "ansibrightgreen": (0x00, 0xFF, 0x00),
+ "ansibrightyellow": (0xFF, 0xFF, 0x00),
+ "ansibrightblue": (0x00, 0x00, 0xFF),
+ "ansibrightmagenta": (0xFF, 0x00, 0xFF),
+ "ansibrightcyan": (0x00, 0xFF, 0xFF),
+}
+
+
+assert set(FG_ANSI_COLORS) == set(ANSI_COLOR_NAMES)
+assert set(BG_ANSI_COLORS) == set(ANSI_COLOR_NAMES)
+assert set(ANSI_COLORS_TO_RGB) == set(ANSI_COLOR_NAMES)
+
+
+def _get_closest_ansi_color(r: int, g: int, b: int, exclude: Sequence[str] = ()) -> str:
+ """
+ Find closest ANSI color. Return it by name.
+
+ :param r: Red (Between 0 and 255.)
+ :param g: Green (Between 0 and 255.)
+ :param b: Blue (Between 0 and 255.)
+ :param exclude: A tuple of color names to exclude. (E.g. ``('ansired', )``.)
+ """
+ exclude = list(exclude)
+
+ # When we have a bit of saturation, avoid the gray-like colors, otherwise,
+ # too often the distance to the gray color is less.
+ saturation = abs(r - g) + abs(g - b) + abs(b - r) # Between 0..510
+
+ if saturation > 30:
+ exclude.extend(["ansilightgray", "ansidarkgray", "ansiwhite", "ansiblack"])
+
+ # Take the closest color.
+ # (Thanks to Pygments for this part.)
+ distance = 257 * 257 * 3 # "infinity" (>distance from #000000 to #ffffff)
+ match = "ansidefault"
+
+ for name, (r2, g2, b2) in ANSI_COLORS_TO_RGB.items():
+ if name != "ansidefault" and name not in exclude:
+ d = (r - r2) ** 2 + (g - g2) ** 2 + (b - b2) ** 2
+
+ if d < distance:
+ match = name
+ distance = d
+
+ return match
+
+
+_ColorCodeAndName = Tuple[int, str]
+
+
+class _16ColorCache:
+ """
+ Cache which maps (r, g, b) tuples to 16 ansi colors.
+
+ :param bg: Cache for background colors, instead of foreground.
+ """
+
+ def __init__(self, bg: bool = False) -> None:
+ self.bg = bg
+ self._cache: dict[Hashable, _ColorCodeAndName] = {}
+
+ def get_code(
+ self, value: tuple[int, int, int], exclude: Sequence[str] = ()
+ ) -> _ColorCodeAndName:
+ """
+ Return a (ansi_code, ansi_name) tuple. (E.g. ``(44, 'ansiblue')``.) for
+ a given (r,g,b) value.
+ """
+ key: Hashable = (value, tuple(exclude))
+ cache = self._cache
+
+ if key not in cache:
+ cache[key] = self._get(value, exclude)
+
+ return cache[key]
+
+ def _get(
+ self, value: tuple[int, int, int], exclude: Sequence[str] = ()
+ ) -> _ColorCodeAndName:
+ r, g, b = value
+ match = _get_closest_ansi_color(r, g, b, exclude=exclude)
+
+ # Turn color name into code.
+ if self.bg:
+ code = BG_ANSI_COLORS[match]
+ else:
+ code = FG_ANSI_COLORS[match]
+
+ return code, match
+
+
+class _256ColorCache(Dict[Tuple[int, int, int], int]):
+ """
+ Cache which maps (r, g, b) tuples to 256 colors.
+ """
+
+ def __init__(self) -> None:
+ # Build color table.
+ colors: list[tuple[int, int, int]] = []
+
+ # colors 0..15: 16 basic colors
+ colors.append((0x00, 0x00, 0x00)) # 0
+ colors.append((0xCD, 0x00, 0x00)) # 1
+ colors.append((0x00, 0xCD, 0x00)) # 2
+ colors.append((0xCD, 0xCD, 0x00)) # 3
+ colors.append((0x00, 0x00, 0xEE)) # 4
+ colors.append((0xCD, 0x00, 0xCD)) # 5
+ colors.append((0x00, 0xCD, 0xCD)) # 6
+ colors.append((0xE5, 0xE5, 0xE5)) # 7
+ colors.append((0x7F, 0x7F, 0x7F)) # 8
+ colors.append((0xFF, 0x00, 0x00)) # 9
+ colors.append((0x00, 0xFF, 0x00)) # 10
+ colors.append((0xFF, 0xFF, 0x00)) # 11
+ colors.append((0x5C, 0x5C, 0xFF)) # 12
+ colors.append((0xFF, 0x00, 0xFF)) # 13
+ colors.append((0x00, 0xFF, 0xFF)) # 14
+ colors.append((0xFF, 0xFF, 0xFF)) # 15
+
+ # colors 16..232: the 6x6x6 color cube
+ valuerange = (0x00, 0x5F, 0x87, 0xAF, 0xD7, 0xFF)
+
+ for i in range(217):
+ r = valuerange[(i // 36) % 6]
+ g = valuerange[(i // 6) % 6]
+ b = valuerange[i % 6]
+ colors.append((r, g, b))
+
+ # colors 233..253: grayscale
+ for i in range(1, 22):
+ v = 8 + i * 10
+ colors.append((v, v, v))
+
+ self.colors = colors
+
+ def __missing__(self, value: tuple[int, int, int]) -> int:
+ r, g, b = value
+
+ # Find closest color.
+ # (Thanks to Pygments for this!)
+ distance = 257 * 257 * 3 # "infinity" (>distance from #000000 to #ffffff)
+ match = 0
+
+ for i, (r2, g2, b2) in enumerate(self.colors):
+ if i >= 16: # XXX: We ignore the 16 ANSI colors when mapping RGB
+ # to the 256 colors, because these highly depend on
+ # the color scheme of the terminal.
+ d = (r - r2) ** 2 + (g - g2) ** 2 + (b - b2) ** 2
+
+ if d < distance:
+ match = i
+ distance = d
+
+ # Turn color name into code.
+ self[value] = match
+ return match
+
+
+_16_fg_colors = _16ColorCache(bg=False)
+_16_bg_colors = _16ColorCache(bg=True)
+_256_colors = _256ColorCache()
+
+
+class _EscapeCodeCache(Dict[Attrs, str]):
+ """
+ Cache for VT100 escape codes. It maps
+ (fgcolor, bgcolor, bold, underline, strike, italic, blink, reverse, hidden, dim) tuples to VT100
+ escape sequences.
+
+ :param true_color: When True, use 24bit colors instead of 256 colors.
+ """
+
+ def __init__(self, color_depth: ColorDepth) -> None:
+ self.color_depth = color_depth
+
+ def __missing__(self, attrs: Attrs) -> str:
+ (
+ fgcolor,
+ bgcolor,
+ bold,
+ underline,
+ strike,
+ italic,
+ blink,
+ reverse,
+ hidden,
+ dim,
+ ) = attrs
+ parts: list[str] = []
+
+ parts.extend(self._colors_to_code(fgcolor or "", bgcolor or ""))
+
+ if bold:
+ parts.append("1")
+ if dim:
+ parts.append("2")
+ if italic:
+ parts.append("3")
+ if blink:
+ parts.append("5")
+ if underline:
+ parts.append("4")
+ if reverse:
+ parts.append("7")
+ if hidden:
+ parts.append("8")
+ if strike:
+ parts.append("9")
+
+ if parts:
+ result = "\x1b[0;" + ";".join(parts) + "m"
+ else:
+ result = "\x1b[0m"
+
+ self[attrs] = result
+ return result
+
+ def _color_name_to_rgb(self, color: str) -> tuple[int, int, int]:
+ "Turn 'ffffff', into (0xff, 0xff, 0xff)."
+ try:
+ rgb = int(color, 16)
+ except ValueError:
+ raise
+ else:
+ r = (rgb >> 16) & 0xFF
+ g = (rgb >> 8) & 0xFF
+ b = rgb & 0xFF
+ return r, g, b
+
+ def _colors_to_code(self, fg_color: str, bg_color: str) -> Iterable[str]:
+ """
+ Return a tuple with the vt100 values that represent this color.
+ """
+ # When requesting ANSI colors only, and both fg/bg color were converted
+ # to ANSI, ensure that the foreground and background color are not the
+ # same. (Unless they were explicitly defined to be the same color.)
+ fg_ansi = ""
+
+ def get(color: str, bg: bool) -> list[int]:
+ nonlocal fg_ansi
+
+ table = BG_ANSI_COLORS if bg else FG_ANSI_COLORS
+
+ if not color or self.color_depth == ColorDepth.DEPTH_1_BIT:
+ return []
+
+ # 16 ANSI colors. (Given by name.)
+ elif color in table:
+ return [table[color]]
+
+ # RGB colors. (Defined as 'ffffff'.)
+ else:
+ try:
+ rgb = self._color_name_to_rgb(color)
+ except ValueError:
+ return []
+
+ # When only 16 colors are supported, use that.
+ if self.color_depth == ColorDepth.DEPTH_4_BIT:
+ if bg: # Background.
+ if fg_color != bg_color:
+ exclude = [fg_ansi]
+ else:
+ exclude = []
+ code, name = _16_bg_colors.get_code(rgb, exclude=exclude)
+ return [code]
+ else: # Foreground.
+ code, name = _16_fg_colors.get_code(rgb)
+ fg_ansi = name
+ return [code]
+
+ # True colors. (Only when this feature is enabled.)
+ elif self.color_depth == ColorDepth.DEPTH_24_BIT:
+ r, g, b = rgb
+ return [(48 if bg else 38), 2, r, g, b]
+
+ # 256 RGB colors.
+ else:
+ return [(48 if bg else 38), 5, _256_colors[rgb]]
+
+ result: list[int] = []
+ result.extend(get(fg_color, False))
+ result.extend(get(bg_color, True))
+
+ return map(str, result)
+
+
+def _get_size(fileno: int) -> tuple[int, int]:
+ """
+ Get the size of this pseudo terminal.
+
+ :param fileno: stdout.fileno()
+ :returns: A (rows, cols) tuple.
+ """
+ size = os.get_terminal_size(fileno)
+ return size.lines, size.columns
+
+
+class Vt100_Output(Output):
+ """
+ :param get_size: A callable which returns the `Size` of the output terminal.
+ :param stdout: Any object with has a `write` and `flush` method + an 'encoding' property.
+ :param term: The terminal environment variable. (xterm, xterm-256color, linux, ...)
+ :param enable_cpr: When `True` (the default), send "cursor position
+ request" escape sequences to the output in order to detect the cursor
+ position. That way, we can properly determine how much space there is
+ available for the UI (especially for drop down menus) to render. The
+ `Renderer` will still try to figure out whether the current terminal
+ does respond to CPR escapes. When `False`, never attempt to send CPR
+ requests.
+ """
+
+ # For the error messages. Only display "Output is not a terminal" once per
+ # file descriptor.
+ _fds_not_a_terminal: set[int] = set()
+
+ def __init__(
+ self,
+ stdout: TextIO,
+ get_size: Callable[[], Size],
+ term: str | None = None,
+ default_color_depth: ColorDepth | None = None,
+ enable_bell: bool = True,
+ enable_cpr: bool = True,
+ ) -> None:
+ assert all(hasattr(stdout, a) for a in ("write", "flush"))
+
+ self._buffer: list[str] = []
+ self.stdout: TextIO = stdout
+ self.default_color_depth = default_color_depth
+ self._get_size = get_size
+ self.term = term
+ self.enable_bell = enable_bell
+ self.enable_cpr = enable_cpr
+
+ # Cache for escape codes.
+ self._escape_code_caches: dict[ColorDepth, _EscapeCodeCache] = {
+ ColorDepth.DEPTH_1_BIT: _EscapeCodeCache(ColorDepth.DEPTH_1_BIT),
+ ColorDepth.DEPTH_4_BIT: _EscapeCodeCache(ColorDepth.DEPTH_4_BIT),
+ ColorDepth.DEPTH_8_BIT: _EscapeCodeCache(ColorDepth.DEPTH_8_BIT),
+ ColorDepth.DEPTH_24_BIT: _EscapeCodeCache(ColorDepth.DEPTH_24_BIT),
+ }
+
+ # Keep track of whether the cursor shape was ever changed.
+ # (We don't restore the cursor shape if it was never changed - by
+ # default, we don't change them.)
+ self._cursor_shape_changed = False
+
+ # Don't hide/show the cursor when this was already done.
+ # (`None` means that we don't know whether the cursor is visible or
+ # not.)
+ self._cursor_visible: bool | None = None
+
+ @classmethod
+ def from_pty(
+ cls,
+ stdout: TextIO,
+ term: str | None = None,
+ default_color_depth: ColorDepth | None = None,
+ enable_bell: bool = True,
+ ) -> Vt100_Output:
+ """
+ Create an Output class from a pseudo terminal.
+ (This will take the dimensions by reading the pseudo
+ terminal attributes.)
+ """
+ fd: int | None
+ # Normally, this requires a real TTY device, but people instantiate
+ # this class often during unit tests as well. For convenience, we print
+ # an error message, use standard dimensions, and go on.
+ try:
+ fd = stdout.fileno()
+ except io.UnsupportedOperation:
+ fd = None
+
+ if not stdout.isatty() and (fd is None or fd not in cls._fds_not_a_terminal):
+ msg = "Warning: Output is not a terminal (fd=%r).\n"
+ sys.stderr.write(msg % fd)
+ sys.stderr.flush()
+ if fd is not None:
+ cls._fds_not_a_terminal.add(fd)
+
+ def get_size() -> Size:
+ # If terminal (incorrectly) reports its size as 0, pick a
+ # reasonable default. See
+ # https://github.com/ipython/ipython/issues/10071
+ rows, columns = (None, None)
+
+ # It is possible that `stdout` is no longer a TTY device at this
+ # point. In that case we get an `OSError` in the ioctl call in
+ # `get_size`. See:
+ # https://github.com/prompt-toolkit/python-prompt-toolkit/pull/1021
+ try:
+ rows, columns = _get_size(stdout.fileno())
+ except OSError:
+ pass
+ return Size(rows=rows or 24, columns=columns or 80)
+
+ return cls(
+ stdout,
+ get_size,
+ term=term,
+ default_color_depth=default_color_depth,
+ enable_bell=enable_bell,
+ )
+
+ def get_size(self) -> Size:
+ return self._get_size()
+
+ def fileno(self) -> int:
+ "Return file descriptor."
+ return self.stdout.fileno()
+
+ def encoding(self) -> str:
+ "Return encoding used for stdout."
+ return self.stdout.encoding
+
+ def write_raw(self, data: str) -> None:
+ """
+ Write raw data to output.
+ """
+ self._buffer.append(data)
+
+ def write(self, data: str) -> None:
+ """
+ Write text to output.
+ (Removes vt100 escape codes. -- used for safely writing text.)
+ """
+ self._buffer.append(data.replace("\x1b", "?"))
+
+ def set_title(self, title: str) -> None:
+ """
+ Set terminal title.
+ """
+ if self.term not in (
+ "linux",
+ "eterm-color",
+ ): # Not supported by the Linux console.
+ self.write_raw(
+ "\x1b]2;{}\x07".format(title.replace("\x1b", "").replace("\x07", ""))
+ )
+
+ def clear_title(self) -> None:
+ self.set_title("")
+
+ def erase_screen(self) -> None:
+ """
+ Erases the screen with the background color and moves the cursor to
+ home.
+ """
+ self.write_raw("\x1b[2J")
+
+ def enter_alternate_screen(self) -> None:
+ self.write_raw("\x1b[?1049h\x1b[H")
+
+ def quit_alternate_screen(self) -> None:
+ self.write_raw("\x1b[?1049l")
+
+ def enable_mouse_support(self) -> None:
+ self.write_raw("\x1b[?1000h")
+
+ # Enable mouse-drag support.
+ self.write_raw("\x1b[?1003h")
+
+ # Enable urxvt Mouse mode. (For terminals that understand this.)
+ self.write_raw("\x1b[?1015h")
+
+ # Also enable Xterm SGR mouse mode. (For terminals that understand this.)
+ self.write_raw("\x1b[?1006h")
+
+ # Note: E.g. lxterminal understands 1000h, but not the urxvt or sgr
+ # extensions.
+
+ def disable_mouse_support(self) -> None:
+ self.write_raw("\x1b[?1000l")
+ self.write_raw("\x1b[?1015l")
+ self.write_raw("\x1b[?1006l")
+ self.write_raw("\x1b[?1003l")
+
+ def erase_end_of_line(self) -> None:
+ """
+ Erases from the current cursor position to the end of the current line.
+ """
+ self.write_raw("\x1b[K")
+
+ def erase_down(self) -> None:
+ """
+ Erases the screen from the current line down to the bottom of the
+ screen.
+ """
+ self.write_raw("\x1b[J")
+
+ def reset_attributes(self) -> None:
+ self.write_raw("\x1b[0m")
+
+ def set_attributes(self, attrs: Attrs, color_depth: ColorDepth) -> None:
+ """
+ Create new style and output.
+
+ :param attrs: `Attrs` instance.
+ """
+ # Get current depth.
+ escape_code_cache = self._escape_code_caches[color_depth]
+
+ # Write escape character.
+ self.write_raw(escape_code_cache[attrs])
+
+ def disable_autowrap(self) -> None:
+ self.write_raw("\x1b[?7l")
+
+ def enable_autowrap(self) -> None:
+ self.write_raw("\x1b[?7h")
+
+ def enable_bracketed_paste(self) -> None:
+ self.write_raw("\x1b[?2004h")
+
+ def disable_bracketed_paste(self) -> None:
+ self.write_raw("\x1b[?2004l")
+
+ def reset_cursor_key_mode(self) -> None:
+ """
+ For vt100 only.
+ Put the terminal in cursor mode (instead of application mode).
+ """
+ # Put the terminal in cursor mode. (Instead of application mode.)
+ self.write_raw("\x1b[?1l")
+
+ def cursor_goto(self, row: int = 0, column: int = 0) -> None:
+ """
+ Move cursor position.
+ """
+ self.write_raw("\x1b[%i;%iH" % (row, column))
+
+ def cursor_up(self, amount: int) -> None:
+ if amount == 0:
+ pass
+ elif amount == 1:
+ self.write_raw("\x1b[A")
+ else:
+ self.write_raw("\x1b[%iA" % amount)
+
+ def cursor_down(self, amount: int) -> None:
+ if amount == 0:
+ pass
+ elif amount == 1:
+ # Note: Not the same as '\n', '\n' can cause the window content to
+ # scroll.
+ self.write_raw("\x1b[B")
+ else:
+ self.write_raw("\x1b[%iB" % amount)
+
+ def cursor_forward(self, amount: int) -> None:
+ if amount == 0:
+ pass
+ elif amount == 1:
+ self.write_raw("\x1b[C")
+ else:
+ self.write_raw("\x1b[%iC" % amount)
+
+ def cursor_backward(self, amount: int) -> None:
+ if amount == 0:
+ pass
+ elif amount == 1:
+ self.write_raw("\b") # '\x1b[D'
+ else:
+ self.write_raw("\x1b[%iD" % amount)
+
+ def hide_cursor(self) -> None:
+ if self._cursor_visible in (True, None):
+ self._cursor_visible = False
+ self.write_raw("\x1b[?25l")
+
+ def show_cursor(self) -> None:
+ if self._cursor_visible in (False, None):
+ self._cursor_visible = True
+ self.write_raw("\x1b[?12l\x1b[?25h") # Stop blinking cursor and show.
+
+ def set_cursor_shape(self, cursor_shape: CursorShape) -> None:
+ if cursor_shape == CursorShape._NEVER_CHANGE:
+ return
+
+ self._cursor_shape_changed = True
+ self.write_raw(
+ {
+ CursorShape.BLOCK: "\x1b[2 q",
+ CursorShape.BEAM: "\x1b[6 q",
+ CursorShape.UNDERLINE: "\x1b[4 q",
+ CursorShape.BLINKING_BLOCK: "\x1b[1 q",
+ CursorShape.BLINKING_BEAM: "\x1b[5 q",
+ CursorShape.BLINKING_UNDERLINE: "\x1b[3 q",
+ }.get(cursor_shape, "")
+ )
+
+ def reset_cursor_shape(self) -> None:
+ "Reset cursor shape."
+ # (Only reset cursor shape, if we ever changed it.)
+ if self._cursor_shape_changed:
+ self._cursor_shape_changed = False
+
+ # Reset cursor shape.
+ self.write_raw("\x1b[0 q")
+
+ def flush(self) -> None:
+ """
+ Write to output stream and flush.
+ """
+ if not self._buffer:
+ return
+
+ data = "".join(self._buffer)
+ self._buffer = []
+
+ flush_stdout(self.stdout, data)
+
+ def ask_for_cpr(self) -> None:
+ """
+ Asks for a cursor position report (CPR).
+ """
+ self.write_raw("\x1b[6n")
+ self.flush()
+
+ @property
+ def responds_to_cpr(self) -> bool:
+ if not self.enable_cpr:
+ return False
+
+ # When the input is a tty, we assume that CPR is supported.
+ # It's not when the input is piped from Pexpect.
+ if os.environ.get("PROMPT_TOOLKIT_NO_CPR", "") == "1":
+ return False
+
+ if is_dumb_terminal(self.term):
+ return False
+ try:
+ return self.stdout.isatty()
+ except ValueError:
+ return False # ValueError: I/O operation on closed file
+
+ def bell(self) -> None:
+ "Sound bell."
+ if self.enable_bell:
+ self.write_raw("\a")
+ self.flush()
+
+ def get_default_color_depth(self) -> ColorDepth:
+ """
+ Return the default color depth for a vt100 terminal, according to the
+ our term value.
+
+ We prefer 256 colors almost always, because this is what most terminals
+ support these days, and is a good default.
+ """
+ if self.default_color_depth is not None:
+ return self.default_color_depth
+
+ term = self.term
+
+ if term is None:
+ return ColorDepth.DEFAULT
+
+ if is_dumb_terminal(term):
+ return ColorDepth.DEPTH_1_BIT
+
+ if term in ("linux", "eterm-color"):
+ return ColorDepth.DEPTH_4_BIT
+
+ return ColorDepth.DEFAULT
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/win32.py b/lib/python3.12/site-packages/prompt_toolkit/output/win32.py
new file mode 100644
index 0000000000000000000000000000000000000000..1c0cc4eec85384f555afd1163ed0922ef86b1129
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/win32.py
@@ -0,0 +1,684 @@
+from __future__ import annotations
+
+import sys
+
+assert sys.platform == "win32"
+
+import os
+from ctypes import ArgumentError, byref, c_char, c_long, c_uint, c_ulong, pointer
+from ctypes.wintypes import DWORD, HANDLE
+from typing import Callable, TextIO, TypeVar
+
+from prompt_toolkit.cursor_shapes import CursorShape
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.styles import ANSI_COLOR_NAMES, Attrs
+from prompt_toolkit.utils import get_cwidth
+from prompt_toolkit.win32_types import (
+ CONSOLE_SCREEN_BUFFER_INFO,
+ COORD,
+ SMALL_RECT,
+ STD_INPUT_HANDLE,
+ STD_OUTPUT_HANDLE,
+)
+
+from ..utils import SPHINX_AUTODOC_RUNNING
+from .base import Output
+from .color_depth import ColorDepth
+
+# Do not import win32-specific stuff when generating documentation.
+# Otherwise RTD would be unable to generate docs for this module.
+if not SPHINX_AUTODOC_RUNNING:
+ from ctypes import windll
+
+
+__all__ = [
+ "Win32Output",
+]
+
+
+def _coord_byval(coord: COORD) -> c_long:
+ """
+ Turns a COORD object into a c_long.
+ This will cause it to be passed by value instead of by reference. (That is what I think at least.)
+
+ When running ``ptipython`` is run (only with IPython), we often got the following error::
+
+ Error in 'SetConsoleCursorPosition'.
+ ArgumentError("argument 2: : wrong type",)
+ argument 2: : wrong type
+
+ It was solved by turning ``COORD`` parameters into a ``c_long`` like this.
+
+ More info: http://msdn.microsoft.com/en-us/library/windows/desktop/ms686025(v=vs.85).aspx
+ """
+ return c_long(coord.Y * 0x10000 | coord.X & 0xFFFF)
+
+
+#: If True: write the output of the renderer also to the following file. This
+#: is very useful for debugging. (e.g.: to see that we don't write more bytes
+#: than required.)
+_DEBUG_RENDER_OUTPUT = False
+_DEBUG_RENDER_OUTPUT_FILENAME = r"prompt-toolkit-windows-output.log"
+
+
+class NoConsoleScreenBufferError(Exception):
+ """
+ Raised when the application is not running inside a Windows Console, but
+ the user tries to instantiate Win32Output.
+ """
+
+ def __init__(self) -> None:
+ # Are we running in 'xterm' on Windows, like git-bash for instance?
+ xterm = "xterm" in os.environ.get("TERM", "")
+
+ if xterm:
+ message = (
+ "Found {}, while expecting a Windows console. "
+ 'Maybe try to run this program using "winpty" '
+ "or run it in cmd.exe instead. Or otherwise, "
+ "in case of Cygwin, use the Python executable "
+ "that is compiled for Cygwin.".format(os.environ["TERM"])
+ )
+ else:
+ message = "No Windows console found. Are you running cmd.exe?"
+ super().__init__(message)
+
+
+_T = TypeVar("_T")
+
+
+class Win32Output(Output):
+ """
+ I/O abstraction for rendering to Windows consoles.
+ (cmd.exe and similar.)
+ """
+
+ def __init__(
+ self,
+ stdout: TextIO,
+ use_complete_width: bool = False,
+ default_color_depth: ColorDepth | None = None,
+ ) -> None:
+ self.use_complete_width = use_complete_width
+ self.default_color_depth = default_color_depth
+
+ self._buffer: list[str] = []
+ self.stdout: TextIO = stdout
+ self.hconsole = HANDLE(windll.kernel32.GetStdHandle(STD_OUTPUT_HANDLE))
+
+ self._in_alternate_screen = False
+ self._hidden = False
+
+ self.color_lookup_table = ColorLookupTable()
+
+ # Remember the default console colors.
+ info = self.get_win32_screen_buffer_info()
+ self.default_attrs = info.wAttributes if info else 15
+
+ if _DEBUG_RENDER_OUTPUT:
+ self.LOG = open(_DEBUG_RENDER_OUTPUT_FILENAME, "ab")
+
+ def fileno(self) -> int:
+ "Return file descriptor."
+ return self.stdout.fileno()
+
+ def encoding(self) -> str:
+ "Return encoding used for stdout."
+ return self.stdout.encoding
+
+ def write(self, data: str) -> None:
+ if self._hidden:
+ data = " " * get_cwidth(data)
+
+ self._buffer.append(data)
+
+ def write_raw(self, data: str) -> None:
+ "For win32, there is no difference between write and write_raw."
+ self.write(data)
+
+ def get_size(self) -> Size:
+ info = self.get_win32_screen_buffer_info()
+
+ # We take the width of the *visible* region as the size. Not the width
+ # of the complete screen buffer. (Unless use_complete_width has been
+ # set.)
+ if self.use_complete_width:
+ width = info.dwSize.X
+ else:
+ width = info.srWindow.Right - info.srWindow.Left
+
+ height = info.srWindow.Bottom - info.srWindow.Top + 1
+
+ # We avoid the right margin, windows will wrap otherwise.
+ maxwidth = info.dwSize.X - 1
+ width = min(maxwidth, width)
+
+ # Create `Size` object.
+ return Size(rows=height, columns=width)
+
+ def _winapi(self, func: Callable[..., _T], *a: object, **kw: object) -> _T:
+ """
+ Flush and call win API function.
+ """
+ self.flush()
+
+ if _DEBUG_RENDER_OUTPUT:
+ self.LOG.write((f"{func.__name__!r}").encode() + b"\n")
+ self.LOG.write(
+ b" " + ", ".join([f"{i!r}" for i in a]).encode("utf-8") + b"\n"
+ )
+ self.LOG.write(
+ b" "
+ + ", ".join([f"{type(i)!r}" for i in a]).encode("utf-8")
+ + b"\n"
+ )
+ self.LOG.flush()
+
+ try:
+ return func(*a, **kw)
+ except ArgumentError as e:
+ if _DEBUG_RENDER_OUTPUT:
+ self.LOG.write((f" Error in {func.__name__!r} {e!r} {e}\n").encode())
+
+ raise
+
+ def get_win32_screen_buffer_info(self) -> CONSOLE_SCREEN_BUFFER_INFO:
+ """
+ Return Screen buffer info.
+ """
+ # NOTE: We don't call the `GetConsoleScreenBufferInfo` API through
+ # `self._winapi`. Doing so causes Python to crash on certain 64bit
+ # Python versions. (Reproduced with 64bit Python 2.7.6, on Windows
+ # 10). It is not clear why. Possibly, it has to do with passing
+ # these objects as an argument, or through *args.
+
+ # The Python documentation contains the following - possibly related - warning:
+ # ctypes does not support passing unions or structures with
+ # bit-fields to functions by value. While this may work on 32-bit
+ # x86, it's not guaranteed by the library to work in the general
+ # case. Unions and structures with bit-fields should always be
+ # passed to functions by pointer.
+
+ # Also see:
+ # - https://github.com/ipython/ipython/issues/10070
+ # - https://github.com/jonathanslenders/python-prompt-toolkit/issues/406
+ # - https://github.com/jonathanslenders/python-prompt-toolkit/issues/86
+
+ self.flush()
+ sbinfo = CONSOLE_SCREEN_BUFFER_INFO()
+ success = windll.kernel32.GetConsoleScreenBufferInfo(
+ self.hconsole, byref(sbinfo)
+ )
+
+ # success = self._winapi(windll.kernel32.GetConsoleScreenBufferInfo,
+ # self.hconsole, byref(sbinfo))
+
+ if success:
+ return sbinfo
+ else:
+ raise NoConsoleScreenBufferError
+
+ def set_title(self, title: str) -> None:
+ """
+ Set terminal title.
+ """
+ self._winapi(windll.kernel32.SetConsoleTitleW, title)
+
+ def clear_title(self) -> None:
+ self._winapi(windll.kernel32.SetConsoleTitleW, "")
+
+ def erase_screen(self) -> None:
+ start = COORD(0, 0)
+ sbinfo = self.get_win32_screen_buffer_info()
+ length = sbinfo.dwSize.X * sbinfo.dwSize.Y
+
+ self.cursor_goto(row=0, column=0)
+ self._erase(start, length)
+
+ def erase_down(self) -> None:
+ sbinfo = self.get_win32_screen_buffer_info()
+ size = sbinfo.dwSize
+
+ start = sbinfo.dwCursorPosition
+ length = (size.X - size.X) + size.X * (size.Y - sbinfo.dwCursorPosition.Y)
+
+ self._erase(start, length)
+
+ def erase_end_of_line(self) -> None:
+ """"""
+ sbinfo = self.get_win32_screen_buffer_info()
+ start = sbinfo.dwCursorPosition
+ length = sbinfo.dwSize.X - sbinfo.dwCursorPosition.X
+
+ self._erase(start, length)
+
+ def _erase(self, start: COORD, length: int) -> None:
+ chars_written = c_ulong()
+
+ self._winapi(
+ windll.kernel32.FillConsoleOutputCharacterA,
+ self.hconsole,
+ c_char(b" "),
+ DWORD(length),
+ _coord_byval(start),
+ byref(chars_written),
+ )
+
+ # Reset attributes.
+ sbinfo = self.get_win32_screen_buffer_info()
+ self._winapi(
+ windll.kernel32.FillConsoleOutputAttribute,
+ self.hconsole,
+ sbinfo.wAttributes,
+ length,
+ _coord_byval(start),
+ byref(chars_written),
+ )
+
+ def reset_attributes(self) -> None:
+ "Reset the console foreground/background color."
+ self._winapi(
+ windll.kernel32.SetConsoleTextAttribute, self.hconsole, self.default_attrs
+ )
+ self._hidden = False
+
+ def set_attributes(self, attrs: Attrs, color_depth: ColorDepth) -> None:
+ (
+ fgcolor,
+ bgcolor,
+ bold,
+ underline,
+ strike,
+ italic,
+ blink,
+ reverse,
+ hidden,
+ dim,
+ ) = attrs
+ self._hidden = bool(hidden)
+
+ # Start from the default attributes.
+ win_attrs: int = self.default_attrs
+
+ if color_depth != ColorDepth.DEPTH_1_BIT:
+ # Override the last four bits: foreground color.
+ if fgcolor:
+ win_attrs = win_attrs & ~0xF
+ win_attrs |= self.color_lookup_table.lookup_fg_color(fgcolor)
+
+ # Override the next four bits: background color.
+ if bgcolor:
+ win_attrs = win_attrs & ~0xF0
+ win_attrs |= self.color_lookup_table.lookup_bg_color(bgcolor)
+
+ # Reverse: swap these four bits groups.
+ if reverse:
+ win_attrs = (
+ (win_attrs & ~0xFF)
+ | ((win_attrs & 0xF) << 4)
+ | ((win_attrs & 0xF0) >> 4)
+ )
+
+ self._winapi(windll.kernel32.SetConsoleTextAttribute, self.hconsole, win_attrs)
+
+ def disable_autowrap(self) -> None:
+ # Not supported by Windows.
+ pass
+
+ def enable_autowrap(self) -> None:
+ # Not supported by Windows.
+ pass
+
+ def cursor_goto(self, row: int = 0, column: int = 0) -> None:
+ pos = COORD(X=column, Y=row)
+ self._winapi(
+ windll.kernel32.SetConsoleCursorPosition, self.hconsole, _coord_byval(pos)
+ )
+
+ def cursor_up(self, amount: int) -> None:
+ sr = self.get_win32_screen_buffer_info().dwCursorPosition
+ pos = COORD(X=sr.X, Y=sr.Y - amount)
+ self._winapi(
+ windll.kernel32.SetConsoleCursorPosition, self.hconsole, _coord_byval(pos)
+ )
+
+ def cursor_down(self, amount: int) -> None:
+ self.cursor_up(-amount)
+
+ def cursor_forward(self, amount: int) -> None:
+ sr = self.get_win32_screen_buffer_info().dwCursorPosition
+ # assert sr.X + amount >= 0, 'Negative cursor position: x=%r amount=%r' % (sr.X, amount)
+
+ pos = COORD(X=max(0, sr.X + amount), Y=sr.Y)
+ self._winapi(
+ windll.kernel32.SetConsoleCursorPosition, self.hconsole, _coord_byval(pos)
+ )
+
+ def cursor_backward(self, amount: int) -> None:
+ self.cursor_forward(-amount)
+
+ def flush(self) -> None:
+ """
+ Write to output stream and flush.
+ """
+ if not self._buffer:
+ # Only flush stdout buffer. (It could be that Python still has
+ # something in its buffer. -- We want to be sure to print that in
+ # the correct color.)
+ self.stdout.flush()
+ return
+
+ data = "".join(self._buffer)
+
+ if _DEBUG_RENDER_OUTPUT:
+ self.LOG.write((f"{data!r}").encode() + b"\n")
+ self.LOG.flush()
+
+ # Print characters one by one. This appears to be the best solution
+ # in order to avoid traces of vertical lines when the completion
+ # menu disappears.
+ for b in data:
+ written = DWORD()
+
+ retval = windll.kernel32.WriteConsoleW(
+ self.hconsole, b, 1, byref(written), None
+ )
+ assert retval != 0
+
+ self._buffer = []
+
+ def get_rows_below_cursor_position(self) -> int:
+ info = self.get_win32_screen_buffer_info()
+ return info.srWindow.Bottom - info.dwCursorPosition.Y + 1
+
+ def scroll_buffer_to_prompt(self) -> None:
+ """
+ To be called before drawing the prompt. This should scroll the console
+ to left, with the cursor at the bottom (if possible).
+ """
+ # Get current window size
+ info = self.get_win32_screen_buffer_info()
+ sr = info.srWindow
+ cursor_pos = info.dwCursorPosition
+
+ result = SMALL_RECT()
+
+ # Scroll to the left.
+ result.Left = 0
+ result.Right = sr.Right - sr.Left
+
+ # Scroll vertical
+ win_height = sr.Bottom - sr.Top
+ if 0 < sr.Bottom - cursor_pos.Y < win_height - 1:
+ # no vertical scroll if cursor already on the screen
+ result.Bottom = sr.Bottom
+ else:
+ result.Bottom = max(win_height, cursor_pos.Y)
+ result.Top = result.Bottom - win_height
+
+ # Scroll API
+ self._winapi(
+ windll.kernel32.SetConsoleWindowInfo, self.hconsole, True, byref(result)
+ )
+
+ def enter_alternate_screen(self) -> None:
+ """
+ Go to alternate screen buffer.
+ """
+ if not self._in_alternate_screen:
+ GENERIC_READ = 0x80000000
+ GENERIC_WRITE = 0x40000000
+
+ # Create a new console buffer and activate that one.
+ handle = HANDLE(
+ self._winapi(
+ windll.kernel32.CreateConsoleScreenBuffer,
+ GENERIC_READ | GENERIC_WRITE,
+ DWORD(0),
+ None,
+ DWORD(1),
+ None,
+ )
+ )
+
+ self._winapi(windll.kernel32.SetConsoleActiveScreenBuffer, handle)
+ self.hconsole = handle
+ self._in_alternate_screen = True
+
+ def quit_alternate_screen(self) -> None:
+ """
+ Make stdout again the active buffer.
+ """
+ if self._in_alternate_screen:
+ stdout = HANDLE(
+ self._winapi(windll.kernel32.GetStdHandle, STD_OUTPUT_HANDLE)
+ )
+ self._winapi(windll.kernel32.SetConsoleActiveScreenBuffer, stdout)
+ self._winapi(windll.kernel32.CloseHandle, self.hconsole)
+ self.hconsole = stdout
+ self._in_alternate_screen = False
+
+ def enable_mouse_support(self) -> None:
+ ENABLE_MOUSE_INPUT = 0x10
+
+ # This `ENABLE_QUICK_EDIT_MODE` flag needs to be cleared for mouse
+ # support to work, but it's possible that it was already cleared
+ # before.
+ ENABLE_QUICK_EDIT_MODE = 0x0040
+
+ handle = HANDLE(windll.kernel32.GetStdHandle(STD_INPUT_HANDLE))
+
+ original_mode = DWORD()
+ self._winapi(windll.kernel32.GetConsoleMode, handle, pointer(original_mode))
+ self._winapi(
+ windll.kernel32.SetConsoleMode,
+ handle,
+ (original_mode.value | ENABLE_MOUSE_INPUT) & ~ENABLE_QUICK_EDIT_MODE,
+ )
+
+ def disable_mouse_support(self) -> None:
+ ENABLE_MOUSE_INPUT = 0x10
+ handle = HANDLE(windll.kernel32.GetStdHandle(STD_INPUT_HANDLE))
+
+ original_mode = DWORD()
+ self._winapi(windll.kernel32.GetConsoleMode, handle, pointer(original_mode))
+ self._winapi(
+ windll.kernel32.SetConsoleMode,
+ handle,
+ original_mode.value & ~ENABLE_MOUSE_INPUT,
+ )
+
+ def hide_cursor(self) -> None:
+ pass
+
+ def show_cursor(self) -> None:
+ pass
+
+ def set_cursor_shape(self, cursor_shape: CursorShape) -> None:
+ pass
+
+ def reset_cursor_shape(self) -> None:
+ pass
+
+ @classmethod
+ def win32_refresh_window(cls) -> None:
+ """
+ Call win32 API to refresh the whole Window.
+
+ This is sometimes necessary when the application paints background
+ for completion menus. When the menu disappears, it leaves traces due
+ to a bug in the Windows Console. Sending a repaint request solves it.
+ """
+ # Get console handle
+ handle = HANDLE(windll.kernel32.GetConsoleWindow())
+
+ RDW_INVALIDATE = 0x0001
+ windll.user32.RedrawWindow(handle, None, None, c_uint(RDW_INVALIDATE))
+
+ def get_default_color_depth(self) -> ColorDepth:
+ """
+ Return the default color depth for a windows terminal.
+
+ Contrary to the Vt100 implementation, this doesn't depend on a $TERM
+ variable.
+ """
+ if self.default_color_depth is not None:
+ return self.default_color_depth
+
+ return ColorDepth.DEPTH_4_BIT
+
+
+class FOREGROUND_COLOR:
+ BLACK = 0x0000
+ BLUE = 0x0001
+ GREEN = 0x0002
+ CYAN = 0x0003
+ RED = 0x0004
+ MAGENTA = 0x0005
+ YELLOW = 0x0006
+ GRAY = 0x0007
+ INTENSITY = 0x0008 # Foreground color is intensified.
+
+
+class BACKGROUND_COLOR:
+ BLACK = 0x0000
+ BLUE = 0x0010
+ GREEN = 0x0020
+ CYAN = 0x0030
+ RED = 0x0040
+ MAGENTA = 0x0050
+ YELLOW = 0x0060
+ GRAY = 0x0070
+ INTENSITY = 0x0080 # Background color is intensified.
+
+
+def _create_ansi_color_dict(
+ color_cls: type[FOREGROUND_COLOR] | type[BACKGROUND_COLOR],
+) -> dict[str, int]:
+ "Create a table that maps the 16 named ansi colors to their Windows code."
+ return {
+ "ansidefault": color_cls.BLACK,
+ "ansiblack": color_cls.BLACK,
+ "ansigray": color_cls.GRAY,
+ "ansibrightblack": color_cls.BLACK | color_cls.INTENSITY,
+ "ansiwhite": color_cls.GRAY | color_cls.INTENSITY,
+ # Low intensity.
+ "ansired": color_cls.RED,
+ "ansigreen": color_cls.GREEN,
+ "ansiyellow": color_cls.YELLOW,
+ "ansiblue": color_cls.BLUE,
+ "ansimagenta": color_cls.MAGENTA,
+ "ansicyan": color_cls.CYAN,
+ # High intensity.
+ "ansibrightred": color_cls.RED | color_cls.INTENSITY,
+ "ansibrightgreen": color_cls.GREEN | color_cls.INTENSITY,
+ "ansibrightyellow": color_cls.YELLOW | color_cls.INTENSITY,
+ "ansibrightblue": color_cls.BLUE | color_cls.INTENSITY,
+ "ansibrightmagenta": color_cls.MAGENTA | color_cls.INTENSITY,
+ "ansibrightcyan": color_cls.CYAN | color_cls.INTENSITY,
+ }
+
+
+FG_ANSI_COLORS = _create_ansi_color_dict(FOREGROUND_COLOR)
+BG_ANSI_COLORS = _create_ansi_color_dict(BACKGROUND_COLOR)
+
+assert set(FG_ANSI_COLORS) == set(ANSI_COLOR_NAMES)
+assert set(BG_ANSI_COLORS) == set(ANSI_COLOR_NAMES)
+
+
+class ColorLookupTable:
+ """
+ Inspired by pygments/formatters/terminal256.py
+ """
+
+ def __init__(self) -> None:
+ self._win32_colors = self._build_color_table()
+
+ # Cache (map color string to foreground and background code).
+ self.best_match: dict[str, tuple[int, int]] = {}
+
+ @staticmethod
+ def _build_color_table() -> list[tuple[int, int, int, int, int]]:
+ """
+ Build an RGB-to-256 color conversion table
+ """
+ FG = FOREGROUND_COLOR
+ BG = BACKGROUND_COLOR
+
+ return [
+ (0x00, 0x00, 0x00, FG.BLACK, BG.BLACK),
+ (0x00, 0x00, 0xAA, FG.BLUE, BG.BLUE),
+ (0x00, 0xAA, 0x00, FG.GREEN, BG.GREEN),
+ (0x00, 0xAA, 0xAA, FG.CYAN, BG.CYAN),
+ (0xAA, 0x00, 0x00, FG.RED, BG.RED),
+ (0xAA, 0x00, 0xAA, FG.MAGENTA, BG.MAGENTA),
+ (0xAA, 0xAA, 0x00, FG.YELLOW, BG.YELLOW),
+ (0x88, 0x88, 0x88, FG.GRAY, BG.GRAY),
+ (0x44, 0x44, 0xFF, FG.BLUE | FG.INTENSITY, BG.BLUE | BG.INTENSITY),
+ (0x44, 0xFF, 0x44, FG.GREEN | FG.INTENSITY, BG.GREEN | BG.INTENSITY),
+ (0x44, 0xFF, 0xFF, FG.CYAN | FG.INTENSITY, BG.CYAN | BG.INTENSITY),
+ (0xFF, 0x44, 0x44, FG.RED | FG.INTENSITY, BG.RED | BG.INTENSITY),
+ (0xFF, 0x44, 0xFF, FG.MAGENTA | FG.INTENSITY, BG.MAGENTA | BG.INTENSITY),
+ (0xFF, 0xFF, 0x44, FG.YELLOW | FG.INTENSITY, BG.YELLOW | BG.INTENSITY),
+ (0x44, 0x44, 0x44, FG.BLACK | FG.INTENSITY, BG.BLACK | BG.INTENSITY),
+ (0xFF, 0xFF, 0xFF, FG.GRAY | FG.INTENSITY, BG.GRAY | BG.INTENSITY),
+ ]
+
+ def _closest_color(self, r: int, g: int, b: int) -> tuple[int, int]:
+ distance = 257 * 257 * 3 # "infinity" (>distance from #000000 to #ffffff)
+ fg_match = 0
+ bg_match = 0
+
+ for r_, g_, b_, fg_, bg_ in self._win32_colors:
+ rd = r - r_
+ gd = g - g_
+ bd = b - b_
+
+ d = rd * rd + gd * gd + bd * bd
+
+ if d < distance:
+ fg_match = fg_
+ bg_match = bg_
+ distance = d
+ return fg_match, bg_match
+
+ def _color_indexes(self, color: str) -> tuple[int, int]:
+ indexes = self.best_match.get(color, None)
+ if indexes is None:
+ try:
+ rgb = int(str(color), 16)
+ except ValueError:
+ rgb = 0
+
+ r = (rgb >> 16) & 0xFF
+ g = (rgb >> 8) & 0xFF
+ b = rgb & 0xFF
+ indexes = self._closest_color(r, g, b)
+ self.best_match[color] = indexes
+ return indexes
+
+ def lookup_fg_color(self, fg_color: str) -> int:
+ """
+ Return the color for use in the
+ `windll.kernel32.SetConsoleTextAttribute` API call.
+
+ :param fg_color: Foreground as text. E.g. 'ffffff' or 'red'
+ """
+ # Foreground.
+ if fg_color in FG_ANSI_COLORS:
+ return FG_ANSI_COLORS[fg_color]
+ else:
+ return self._color_indexes(fg_color)[0]
+
+ def lookup_bg_color(self, bg_color: str) -> int:
+ """
+ Return the color for use in the
+ `windll.kernel32.SetConsoleTextAttribute` API call.
+
+ :param bg_color: Background as text. E.g. 'ffffff' or 'red'
+ """
+ # Background.
+ if bg_color in BG_ANSI_COLORS:
+ return BG_ANSI_COLORS[bg_color]
+ else:
+ return self._color_indexes(bg_color)[1]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/output/windows10.py b/lib/python3.12/site-packages/prompt_toolkit/output/windows10.py
new file mode 100644
index 0000000000000000000000000000000000000000..2b7e596e468a63b8a763716a2472560466faf74a
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/output/windows10.py
@@ -0,0 +1,133 @@
+from __future__ import annotations
+
+import sys
+
+assert sys.platform == "win32"
+
+from ctypes import byref, windll
+from ctypes.wintypes import DWORD, HANDLE
+from typing import Any, TextIO
+
+from prompt_toolkit.data_structures import Size
+from prompt_toolkit.win32_types import STD_OUTPUT_HANDLE
+
+from .base import Output
+from .color_depth import ColorDepth
+from .vt100 import Vt100_Output
+from .win32 import Win32Output
+
+__all__ = [
+ "Windows10_Output",
+]
+
+# See: https://msdn.microsoft.com/pl-pl/library/windows/desktop/ms686033(v=vs.85).aspx
+ENABLE_PROCESSED_INPUT = 0x0001
+ENABLE_VIRTUAL_TERMINAL_PROCESSING = 0x0004
+
+
+class Windows10_Output:
+ """
+ Windows 10 output abstraction. This enables and uses vt100 escape sequences.
+ """
+
+ def __init__(
+ self, stdout: TextIO, default_color_depth: ColorDepth | None = None
+ ) -> None:
+ self.default_color_depth = default_color_depth
+ self.win32_output = Win32Output(stdout, default_color_depth=default_color_depth)
+ self.vt100_output = Vt100_Output(
+ stdout, lambda: Size(0, 0), default_color_depth=default_color_depth
+ )
+ self._hconsole = HANDLE(windll.kernel32.GetStdHandle(STD_OUTPUT_HANDLE))
+
+ def flush(self) -> None:
+ """
+ Write to output stream and flush.
+ """
+ original_mode = DWORD(0)
+
+ # Remember the previous console mode.
+ windll.kernel32.GetConsoleMode(self._hconsole, byref(original_mode))
+
+ # Enable processing of vt100 sequences.
+ windll.kernel32.SetConsoleMode(
+ self._hconsole,
+ DWORD(ENABLE_PROCESSED_INPUT | ENABLE_VIRTUAL_TERMINAL_PROCESSING),
+ )
+
+ try:
+ self.vt100_output.flush()
+ finally:
+ # Restore console mode.
+ windll.kernel32.SetConsoleMode(self._hconsole, original_mode)
+
+ @property
+ def responds_to_cpr(self) -> bool:
+ return False # We don't need this on Windows.
+
+ def __getattr__(self, name: str) -> Any:
+ # NOTE: Now that we use "virtual terminal input" on
+ # Windows, both input and output are done through
+ # ANSI escape sequences on Windows. This means, we
+ # should enable bracketed paste like on Linux, and
+ # enable mouse support by calling the vt100_output.
+ if name in (
+ "get_size",
+ "get_rows_below_cursor_position",
+ "scroll_buffer_to_prompt",
+ "get_win32_screen_buffer_info",
+ # "enable_mouse_support",
+ # "disable_mouse_support",
+ # "enable_bracketed_paste",
+ # "disable_bracketed_paste",
+ ):
+ return getattr(self.win32_output, name)
+ else:
+ return getattr(self.vt100_output, name)
+
+ def get_default_color_depth(self) -> ColorDepth:
+ """
+ Return the default color depth for a windows terminal.
+
+ Contrary to the Vt100 implementation, this doesn't depend on a $TERM
+ variable.
+ """
+ if self.default_color_depth is not None:
+ return self.default_color_depth
+
+ # Previously, we used `DEPTH_4_BIT`, even on Windows 10. This was
+ # because true color support was added after "Console Virtual Terminal
+ # Sequences" support was added, and there was no good way to detect
+ # what support was given.
+ # 24bit color support was added in 2016, so let's assume it's safe to
+ # take that as a default:
+ # https://devblogs.microsoft.com/commandline/24-bit-color-in-the-windows-console/
+ return ColorDepth.TRUE_COLOR
+
+
+Output.register(Windows10_Output)
+
+
+def is_win_vt100_enabled() -> bool:
+ """
+ Returns True when we're running Windows and VT100 escape sequences are
+ supported.
+ """
+ if sys.platform != "win32":
+ return False
+
+ hconsole = HANDLE(windll.kernel32.GetStdHandle(STD_OUTPUT_HANDLE))
+
+ # Get original console mode.
+ original_mode = DWORD(0)
+ windll.kernel32.GetConsoleMode(hconsole, byref(original_mode))
+
+ try:
+ # Try to enable VT100 sequences.
+ result: int = windll.kernel32.SetConsoleMode(
+ hconsole, DWORD(ENABLE_PROCESSED_INPUT | ENABLE_VIRTUAL_TERMINAL_PROCESSING)
+ )
+
+ return result == 1
+ finally:
+ windll.kernel32.SetConsoleMode(hconsole, original_mode)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/patch_stdout.py b/lib/python3.12/site-packages/prompt_toolkit/patch_stdout.py
new file mode 100644
index 0000000000000000000000000000000000000000..e1f2a7a2c3a582101f290116b729d3de1b73daca
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/patch_stdout.py
@@ -0,0 +1,297 @@
+"""
+patch_stdout
+============
+
+This implements a context manager that ensures that print statements within
+it won't destroy the user interface. The context manager will replace
+`sys.stdout` by something that draws the output above the current prompt,
+rather than overwriting the UI.
+
+Usage::
+
+ with patch_stdout(application):
+ ...
+ application.run()
+ ...
+
+Multiple applications can run in the body of the context manager, one after the
+other.
+"""
+
+from __future__ import annotations
+
+import asyncio
+import queue
+import sys
+import threading
+import time
+from contextlib import contextmanager
+from typing import Generator, TextIO, cast
+
+from .application import get_app_session, run_in_terminal
+from .output import Output
+
+__all__ = [
+ "patch_stdout",
+ "StdoutProxy",
+]
+
+
+@contextmanager
+def patch_stdout(raw: bool = False) -> Generator[None, None, None]:
+ """
+ Replace `sys.stdout` by an :class:`_StdoutProxy` instance.
+
+ Writing to this proxy will make sure that the text appears above the
+ prompt, and that it doesn't destroy the output from the renderer. If no
+ application is curring, the behavior should be identical to writing to
+ `sys.stdout` directly.
+
+ Warning: If a new event loop is installed using `asyncio.set_event_loop()`,
+ then make sure that the context manager is applied after the event loop
+ is changed. Printing to stdout will be scheduled in the event loop
+ that's active when the context manager is created.
+
+ :param raw: (`bool`) When True, vt100 terminal escape sequences are not
+ removed/escaped.
+ """
+ with StdoutProxy(raw=raw) as proxy:
+ original_stdout = sys.stdout
+ original_stderr = sys.stderr
+
+ # Enter.
+ sys.stdout = cast(TextIO, proxy)
+ sys.stderr = cast(TextIO, proxy)
+
+ try:
+ yield
+ finally:
+ sys.stdout = original_stdout
+ sys.stderr = original_stderr
+
+
+class _Done:
+ "Sentinel value for stopping the stdout proxy."
+
+
+class StdoutProxy:
+ """
+ File-like object, which prints everything written to it, output above the
+ current application/prompt. This class is compatible with other file
+ objects and can be used as a drop-in replacement for `sys.stdout` or can
+ for instance be passed to `logging.StreamHandler`.
+
+ The current application, above which we print, is determined by looking
+ what application currently runs in the `AppSession` that is active during
+ the creation of this instance.
+
+ This class can be used as a context manager.
+
+ In order to avoid having to repaint the prompt continuously for every
+ little write, a short delay of `sleep_between_writes` seconds will be added
+ between writes in order to bundle many smaller writes in a short timespan.
+ """
+
+ def __init__(
+ self,
+ sleep_between_writes: float = 0.2,
+ raw: bool = False,
+ ) -> None:
+ self.sleep_between_writes = sleep_between_writes
+ self.raw = raw
+
+ self._lock = threading.RLock()
+ self._buffer: list[str] = []
+
+ # Keep track of the curret app session.
+ self.app_session = get_app_session()
+
+ # See what output is active *right now*. We should do it at this point,
+ # before this `StdoutProxy` instance is possibly assigned to `sys.stdout`.
+ # Otherwise, if `patch_stdout` is used, and no `Output` instance has
+ # been created, then the default output creation code will see this
+ # proxy object as `sys.stdout`, and get in a recursive loop trying to
+ # access `StdoutProxy.isatty()` which will again retrieve the output.
+ self._output: Output = self.app_session.output
+
+ # Flush thread
+ self._flush_queue: queue.Queue[str | _Done] = queue.Queue()
+ self._flush_thread = self._start_write_thread()
+ self.closed = False
+
+ def __enter__(self) -> StdoutProxy:
+ return self
+
+ def __exit__(self, *args: object) -> None:
+ self.close()
+
+ def close(self) -> None:
+ """
+ Stop `StdoutProxy` proxy.
+
+ This will terminate the write thread, make sure everything is flushed
+ and wait for the write thread to finish.
+ """
+ if not self.closed:
+ self._flush_queue.put(_Done())
+ self._flush_thread.join()
+ self.closed = True
+
+ def _start_write_thread(self) -> threading.Thread:
+ thread = threading.Thread(
+ target=self._write_thread,
+ name="patch-stdout-flush-thread",
+ daemon=True,
+ )
+ thread.start()
+ return thread
+
+ def _write_thread(self) -> None:
+ done = False
+
+ while not done:
+ item = self._flush_queue.get()
+
+ if isinstance(item, _Done):
+ break
+
+ # Don't bother calling when we got an empty string.
+ if not item:
+ continue
+
+ text = []
+ text.append(item)
+
+ # Read the rest of the queue if more data was queued up.
+ while True:
+ try:
+ item = self._flush_queue.get_nowait()
+ except queue.Empty:
+ break
+ else:
+ if isinstance(item, _Done):
+ done = True
+ else:
+ text.append(item)
+
+ app_loop = self._get_app_loop()
+ self._write_and_flush(app_loop, "".join(text))
+
+ # If an application was running that requires repainting, then wait
+ # for a very short time, in order to bundle actual writes and avoid
+ # having to repaint to often.
+ if app_loop is not None:
+ time.sleep(self.sleep_between_writes)
+
+ def _get_app_loop(self) -> asyncio.AbstractEventLoop | None:
+ """
+ Return the event loop for the application currently running in our
+ `AppSession`.
+ """
+ app = self.app_session.app
+
+ if app is None:
+ return None
+
+ return app.loop
+
+ def _write_and_flush(
+ self, loop: asyncio.AbstractEventLoop | None, text: str
+ ) -> None:
+ """
+ Write the given text to stdout and flush.
+ If an application is running, use `run_in_terminal`.
+ """
+
+ def write_and_flush() -> None:
+ # Ensure that autowrap is enabled before calling `write`.
+ # XXX: On Windows, the `Windows10_Output` enables/disables VT
+ # terminal processing for every flush. It turns out that this
+ # causes autowrap to be reset (disabled) after each flush. So,
+ # we have to enable it again before writing text.
+ self._output.enable_autowrap()
+
+ if self.raw:
+ self._output.write_raw(text)
+ else:
+ self._output.write(text)
+
+ self._output.flush()
+
+ def write_and_flush_in_loop() -> None:
+ # If an application is running, use `run_in_terminal`, otherwise
+ # call it directly.
+ run_in_terminal(write_and_flush, in_executor=False)
+
+ if loop is None:
+ # No loop, write immediately.
+ write_and_flush()
+ else:
+ # Make sure `write_and_flush` is executed *in* the event loop, not
+ # in another thread.
+ loop.call_soon_threadsafe(write_and_flush_in_loop)
+
+ def _write(self, data: str) -> None:
+ """
+ Note: print()-statements cause to multiple write calls.
+ (write('line') and write('\n')). Of course we don't want to call
+ `run_in_terminal` for every individual call, because that's too
+ expensive, and as long as the newline hasn't been written, the
+ text itself is again overwritten by the rendering of the input
+ command line. Therefor, we have a little buffer which holds the
+ text until a newline is written to stdout.
+ """
+ if "\n" in data:
+ # When there is a newline in the data, write everything before the
+ # newline, including the newline itself.
+ before, after = data.rsplit("\n", 1)
+ to_write = self._buffer + [before, "\n"]
+ self._buffer = [after]
+
+ text = "".join(to_write)
+ self._flush_queue.put(text)
+ else:
+ # Otherwise, cache in buffer.
+ self._buffer.append(data)
+
+ def _flush(self) -> None:
+ text = "".join(self._buffer)
+ self._buffer = []
+ self._flush_queue.put(text)
+
+ def write(self, data: str) -> int:
+ with self._lock:
+ self._write(data)
+
+ return len(data) # Pretend everything was written.
+
+ def flush(self) -> None:
+ """
+ Flush buffered output.
+ """
+ with self._lock:
+ self._flush()
+
+ @property
+ def original_stdout(self) -> TextIO | None:
+ return self._output.stdout or sys.__stdout__
+
+ # Attributes for compatibility with sys.__stdout__:
+
+ def fileno(self) -> int:
+ return self._output.fileno()
+
+ def isatty(self) -> bool:
+ stdout = self._output.stdout
+ if stdout is None:
+ return False
+
+ return stdout.isatty()
+
+ @property
+ def encoding(self) -> str:
+ return self._output.encoding()
+
+ @property
+ def errors(self) -> str:
+ return "strict"
diff --git a/lib/python3.12/site-packages/prompt_toolkit/py.typed b/lib/python3.12/site-packages/prompt_toolkit/py.typed
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/lib/python3.12/site-packages/prompt_toolkit/renderer.py b/lib/python3.12/site-packages/prompt_toolkit/renderer.py
new file mode 100644
index 0000000000000000000000000000000000000000..8d5e03c19e8948531bfd7e1994fefce0da75b8e7
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/renderer.py
@@ -0,0 +1,820 @@
+"""
+Renders the command line on the console.
+(Redraws parts of the input line that were changed.)
+"""
+
+from __future__ import annotations
+
+from asyncio import FIRST_COMPLETED, Future, ensure_future, sleep, wait
+from collections import deque
+from enum import Enum
+from typing import TYPE_CHECKING, Any, Callable, Dict, Hashable
+
+from prompt_toolkit.application.current import get_app
+from prompt_toolkit.cursor_shapes import CursorShape
+from prompt_toolkit.data_structures import Point, Size
+from prompt_toolkit.filters import FilterOrBool, to_filter
+from prompt_toolkit.formatted_text import AnyFormattedText, to_formatted_text
+from prompt_toolkit.layout.mouse_handlers import MouseHandlers
+from prompt_toolkit.layout.screen import Char, Screen, WritePosition
+from prompt_toolkit.output import ColorDepth, Output
+from prompt_toolkit.styles import (
+ Attrs,
+ BaseStyle,
+ DummyStyleTransformation,
+ StyleTransformation,
+)
+
+if TYPE_CHECKING:
+ from prompt_toolkit.application import Application
+ from prompt_toolkit.layout.layout import Layout
+
+
+__all__ = [
+ "Renderer",
+ "print_formatted_text",
+]
+
+
+def _output_screen_diff(
+ app: Application[Any],
+ output: Output,
+ screen: Screen,
+ current_pos: Point,
+ color_depth: ColorDepth,
+ previous_screen: Screen | None,
+ last_style: str | None,
+ is_done: bool, # XXX: drop is_done
+ full_screen: bool,
+ attrs_for_style_string: _StyleStringToAttrsCache,
+ style_string_has_style: _StyleStringHasStyleCache,
+ size: Size,
+ previous_width: int,
+) -> tuple[Point, str | None]:
+ """
+ Render the diff between this screen and the previous screen.
+
+ This takes two `Screen` instances. The one that represents the output like
+ it was during the last rendering and one that represents the current
+ output raster. Looking at these two `Screen` instances, this function will
+ render the difference by calling the appropriate methods of the `Output`
+ object that only paint the changes to the terminal.
+
+ This is some performance-critical code which is heavily optimized.
+ Don't change things without profiling first.
+
+ :param current_pos: Current cursor position.
+ :param last_style: The style string, used for drawing the last drawn
+ character. (Color/attributes.)
+ :param attrs_for_style_string: :class:`._StyleStringToAttrsCache` instance.
+ :param width: The width of the terminal.
+ :param previous_width: The width of the terminal during the last rendering.
+ """
+ width, height = size.columns, size.rows
+
+ #: Variable for capturing the output.
+ write = output.write
+ write_raw = output.write_raw
+
+ # Create locals for the most used output methods.
+ # (Save expensive attribute lookups.)
+ _output_set_attributes = output.set_attributes
+ _output_reset_attributes = output.reset_attributes
+ _output_cursor_forward = output.cursor_forward
+ _output_cursor_up = output.cursor_up
+ _output_cursor_backward = output.cursor_backward
+
+ # Hide cursor before rendering. (Avoid flickering.)
+ output.hide_cursor()
+
+ def reset_attributes() -> None:
+ "Wrapper around Output.reset_attributes."
+ nonlocal last_style
+ _output_reset_attributes()
+ last_style = None # Forget last char after resetting attributes.
+
+ def move_cursor(new: Point) -> Point:
+ "Move cursor to this `new` point. Returns the given Point."
+ current_x, current_y = current_pos.x, current_pos.y
+
+ if new.y > current_y:
+ # Use newlines instead of CURSOR_DOWN, because this might add new lines.
+ # CURSOR_DOWN will never create new lines at the bottom.
+ # Also reset attributes, otherwise the newline could draw a
+ # background color.
+ reset_attributes()
+ write("\r\n" * (new.y - current_y))
+ current_x = 0
+ _output_cursor_forward(new.x)
+ return new
+ elif new.y < current_y:
+ _output_cursor_up(current_y - new.y)
+
+ if current_x >= width - 1:
+ write("\r")
+ _output_cursor_forward(new.x)
+ elif new.x < current_x or current_x >= width - 1:
+ _output_cursor_backward(current_x - new.x)
+ elif new.x > current_x:
+ _output_cursor_forward(new.x - current_x)
+
+ return new
+
+ def output_char(char: Char) -> None:
+ """
+ Write the output of this character.
+ """
+ nonlocal last_style
+
+ # If the last printed character has the same style, don't output the
+ # style again.
+ if last_style == char.style:
+ write(char.char)
+ else:
+ # Look up `Attr` for this style string. Only set attributes if different.
+ # (Two style strings can still have the same formatting.)
+ # Note that an empty style string can have formatting that needs to
+ # be applied, because of style transformations.
+ new_attrs = attrs_for_style_string[char.style]
+ if not last_style or new_attrs != attrs_for_style_string[last_style]:
+ _output_set_attributes(new_attrs, color_depth)
+
+ write(char.char)
+ last_style = char.style
+
+ def get_max_column_index(row: dict[int, Char]) -> int:
+ """
+ Return max used column index, ignoring whitespace (without style) at
+ the end of the line. This is important for people that copy/paste
+ terminal output.
+
+ There are two reasons we are sometimes seeing whitespace at the end:
+ - `BufferControl` adds a trailing space to each line, because it's a
+ possible cursor position, so that the line wrapping won't change if
+ the cursor position moves around.
+ - The `Window` adds a style class to the current line for highlighting
+ (cursor-line).
+ """
+ numbers = (
+ index
+ for index, cell in row.items()
+ if cell.char != " " or style_string_has_style[cell.style]
+ )
+ return max(numbers, default=0)
+
+ # Render for the first time: reset styling.
+ if not previous_screen:
+ reset_attributes()
+
+ # Disable autowrap. (When entering a the alternate screen, or anytime when
+ # we have a prompt. - In the case of a REPL, like IPython, people can have
+ # background threads, and it's hard for debugging if their output is not
+ # wrapped.)
+ if not previous_screen or not full_screen:
+ output.disable_autowrap()
+
+ # When the previous screen has a different size, redraw everything anyway.
+ # Also when we are done. (We might take up less rows, so clearing is important.)
+ if (
+ is_done or not previous_screen or previous_width != width
+ ): # XXX: also consider height??
+ current_pos = move_cursor(Point(x=0, y=0))
+ reset_attributes()
+ output.erase_down()
+
+ previous_screen = Screen()
+
+ # Get height of the screen.
+ # (height changes as we loop over data_buffer, so remember the current value.)
+ # (Also make sure to clip the height to the size of the output.)
+ current_height = min(screen.height, height)
+
+ # Loop over the rows.
+ row_count = min(max(screen.height, previous_screen.height), height)
+
+ for y in range(row_count):
+ new_row = screen.data_buffer[y]
+ previous_row = previous_screen.data_buffer[y]
+ zero_width_escapes_row = screen.zero_width_escapes[y]
+
+ new_max_line_len = min(width - 1, get_max_column_index(new_row))
+ previous_max_line_len = min(width - 1, get_max_column_index(previous_row))
+
+ # Loop over the columns.
+ c = 0 # Column counter.
+ while c <= new_max_line_len:
+ new_char = new_row[c]
+ old_char = previous_row[c]
+ char_width = new_char.width or 1
+
+ # When the old and new character at this position are different,
+ # draw the output. (Because of the performance, we don't call
+ # `Char.__ne__`, but inline the same expression.)
+ if new_char.char != old_char.char or new_char.style != old_char.style:
+ current_pos = move_cursor(Point(x=c, y=y))
+
+ # Send injected escape sequences to output.
+ if c in zero_width_escapes_row:
+ write_raw(zero_width_escapes_row[c])
+
+ output_char(new_char)
+ current_pos = Point(x=current_pos.x + char_width, y=current_pos.y)
+
+ c += char_width
+
+ # If the new line is shorter, trim it.
+ if previous_screen and new_max_line_len < previous_max_line_len:
+ current_pos = move_cursor(Point(x=new_max_line_len + 1, y=y))
+ reset_attributes()
+ output.erase_end_of_line()
+
+ # Correctly reserve vertical space as required by the layout.
+ # When this is a new screen (drawn for the first time), or for some reason
+ # higher than the previous one. Move the cursor once to the bottom of the
+ # output. That way, we're sure that the terminal scrolls up, even when the
+ # lower lines of the canvas just contain whitespace.
+
+ # The most obvious reason that we actually want this behavior is the avoid
+ # the artifact of the input scrolling when the completion menu is shown.
+ # (If the scrolling is actually wanted, the layout can still be build in a
+ # way to behave that way by setting a dynamic height.)
+ if current_height > previous_screen.height:
+ current_pos = move_cursor(Point(x=0, y=current_height - 1))
+
+ # Move cursor:
+ if is_done:
+ current_pos = move_cursor(Point(x=0, y=current_height))
+ output.erase_down()
+ else:
+ current_pos = move_cursor(screen.get_cursor_position(app.layout.current_window))
+
+ if is_done or not full_screen:
+ output.enable_autowrap()
+
+ # Always reset the color attributes. This is important because a background
+ # thread could print data to stdout and we want that to be displayed in the
+ # default colors. (Also, if a background color has been set, many terminals
+ # give weird artifacts on resize events.)
+ reset_attributes()
+
+ if screen.show_cursor:
+ output.show_cursor()
+
+ return current_pos, last_style
+
+
+class HeightIsUnknownError(Exception):
+ "Information unavailable. Did not yet receive the CPR response."
+
+
+class _StyleStringToAttrsCache(Dict[str, Attrs]):
+ """
+ A cache structure that maps style strings to :class:`.Attr`.
+ (This is an important speed up.)
+ """
+
+ def __init__(
+ self,
+ get_attrs_for_style_str: Callable[[str], Attrs],
+ style_transformation: StyleTransformation,
+ ) -> None:
+ self.get_attrs_for_style_str = get_attrs_for_style_str
+ self.style_transformation = style_transformation
+
+ def __missing__(self, style_str: str) -> Attrs:
+ attrs = self.get_attrs_for_style_str(style_str)
+ attrs = self.style_transformation.transform_attrs(attrs)
+
+ self[style_str] = attrs
+ return attrs
+
+
+class _StyleStringHasStyleCache(Dict[str, bool]):
+ """
+ Cache for remember which style strings don't render the default output
+ style (default fg/bg, no underline and no reverse and no blink). That way
+ we know that we should render these cells, even when they're empty (when
+ they contain a space).
+
+ Note: we don't consider bold/italic/hidden because they don't change the
+ output if there's no text in the cell.
+ """
+
+ def __init__(self, style_string_to_attrs: dict[str, Attrs]) -> None:
+ self.style_string_to_attrs = style_string_to_attrs
+
+ def __missing__(self, style_str: str) -> bool:
+ attrs = self.style_string_to_attrs[style_str]
+ is_default = bool(
+ attrs.color
+ or attrs.bgcolor
+ or attrs.underline
+ or attrs.strike
+ or attrs.blink
+ or attrs.reverse
+ )
+
+ self[style_str] = is_default
+ return is_default
+
+
+class CPR_Support(Enum):
+ "Enum: whether or not CPR is supported."
+
+ SUPPORTED = "SUPPORTED"
+ NOT_SUPPORTED = "NOT_SUPPORTED"
+ UNKNOWN = "UNKNOWN"
+
+
+class Renderer:
+ """
+ Typical usage:
+
+ ::
+
+ output = Vt100_Output.from_pty(sys.stdout)
+ r = Renderer(style, output)
+ r.render(app, layout=...)
+ """
+
+ CPR_TIMEOUT = 2 # Time to wait until we consider CPR to be not supported.
+
+ def __init__(
+ self,
+ style: BaseStyle,
+ output: Output,
+ full_screen: bool = False,
+ mouse_support: FilterOrBool = False,
+ cpr_not_supported_callback: Callable[[], None] | None = None,
+ ) -> None:
+ self.style = style
+ self.output = output
+ self.full_screen = full_screen
+ self.mouse_support = to_filter(mouse_support)
+ self.cpr_not_supported_callback = cpr_not_supported_callback
+
+ # TODO: Move following state flags into `Vt100_Output`, similar to
+ # `_cursor_shape_changed` and `_cursor_visible`. But then also
+ # adjust the `Win32Output` to not call win32 APIs if nothing has
+ # to be changed.
+
+ self._in_alternate_screen = False
+ self._mouse_support_enabled = False
+ self._bracketed_paste_enabled = False
+ self._cursor_key_mode_reset = False
+
+ # Future set when we are waiting for a CPR flag.
+ self._waiting_for_cpr_futures: deque[Future[None]] = deque()
+ self.cpr_support = CPR_Support.UNKNOWN
+
+ if not output.responds_to_cpr:
+ self.cpr_support = CPR_Support.NOT_SUPPORTED
+
+ # Cache for the style.
+ self._attrs_for_style: _StyleStringToAttrsCache | None = None
+ self._style_string_has_style: _StyleStringHasStyleCache | None = None
+ self._last_style_hash: Hashable | None = None
+ self._last_transformation_hash: Hashable | None = None
+ self._last_color_depth: ColorDepth | None = None
+
+ self.reset(_scroll=True)
+
+ def reset(self, _scroll: bool = False, leave_alternate_screen: bool = True) -> None:
+ # Reset position
+ self._cursor_pos = Point(x=0, y=0)
+
+ # Remember the last screen instance between renderers. This way,
+ # we can create a `diff` between two screens and only output the
+ # difference. It's also to remember the last height. (To show for
+ # instance a toolbar at the bottom position.)
+ self._last_screen: Screen | None = None
+ self._last_size: Size | None = None
+ self._last_style: str | None = None
+ self._last_cursor_shape: CursorShape | None = None
+
+ # Default MouseHandlers. (Just empty.)
+ self.mouse_handlers = MouseHandlers()
+
+ #: Space from the top of the layout, until the bottom of the terminal.
+ #: We don't know this until a `report_absolute_cursor_row` call.
+ self._min_available_height = 0
+
+ # In case of Windows, also make sure to scroll to the current cursor
+ # position. (Only when rendering the first time.)
+ # It does nothing for vt100 terminals.
+ if _scroll:
+ self.output.scroll_buffer_to_prompt()
+
+ # Quit alternate screen.
+ if self._in_alternate_screen and leave_alternate_screen:
+ self.output.quit_alternate_screen()
+ self._in_alternate_screen = False
+
+ # Disable mouse support.
+ if self._mouse_support_enabled:
+ self.output.disable_mouse_support()
+ self._mouse_support_enabled = False
+
+ # Disable bracketed paste.
+ if self._bracketed_paste_enabled:
+ self.output.disable_bracketed_paste()
+ self._bracketed_paste_enabled = False
+
+ self.output.reset_cursor_shape()
+ self.output.show_cursor()
+
+ # NOTE: No need to set/reset cursor key mode here.
+
+ # Flush output. `disable_mouse_support` needs to write to stdout.
+ self.output.flush()
+
+ @property
+ def last_rendered_screen(self) -> Screen | None:
+ """
+ The `Screen` class that was generated during the last rendering.
+ This can be `None`.
+ """
+ return self._last_screen
+
+ @property
+ def height_is_known(self) -> bool:
+ """
+ True when the height from the cursor until the bottom of the terminal
+ is known. (It's often nicer to draw bottom toolbars only if the height
+ is known, in order to avoid flickering when the CPR response arrives.)
+ """
+ if self.full_screen or self._min_available_height > 0:
+ return True
+ try:
+ self._min_available_height = self.output.get_rows_below_cursor_position()
+ return True
+ except NotImplementedError:
+ return False
+
+ @property
+ def rows_above_layout(self) -> int:
+ """
+ Return the number of rows visible in the terminal above the layout.
+ """
+ if self._in_alternate_screen:
+ return 0
+ elif self._min_available_height > 0:
+ total_rows = self.output.get_size().rows
+ last_screen_height = self._last_screen.height if self._last_screen else 0
+ return total_rows - max(self._min_available_height, last_screen_height)
+ else:
+ raise HeightIsUnknownError("Rows above layout is unknown.")
+
+ def request_absolute_cursor_position(self) -> None:
+ """
+ Get current cursor position.
+
+ We do this to calculate the minimum available height that we can
+ consume for rendering the prompt. This is the available space below te
+ cursor.
+
+ For vt100: Do CPR request. (answer will arrive later.)
+ For win32: Do API call. (Answer comes immediately.)
+ """
+ # Only do this request when the cursor is at the top row. (after a
+ # clear or reset). We will rely on that in `report_absolute_cursor_row`.
+ assert self._cursor_pos.y == 0
+
+ # In full-screen mode, always use the total height as min-available-height.
+ if self.full_screen:
+ self._min_available_height = self.output.get_size().rows
+ return
+
+ # For Win32, we have an API call to get the number of rows below the
+ # cursor.
+ try:
+ self._min_available_height = self.output.get_rows_below_cursor_position()
+ return
+ except NotImplementedError:
+ pass
+
+ # Use CPR.
+ if self.cpr_support == CPR_Support.NOT_SUPPORTED:
+ return
+
+ def do_cpr() -> None:
+ # Asks for a cursor position report (CPR).
+ self._waiting_for_cpr_futures.append(Future())
+ self.output.ask_for_cpr()
+
+ if self.cpr_support == CPR_Support.SUPPORTED:
+ do_cpr()
+ return
+
+ # If we don't know whether CPR is supported, only do a request if
+ # none is pending, and test it, using a timer.
+ if self.waiting_for_cpr:
+ return
+
+ do_cpr()
+
+ async def timer() -> None:
+ await sleep(self.CPR_TIMEOUT)
+
+ # Not set in the meantime -> not supported.
+ if self.cpr_support == CPR_Support.UNKNOWN:
+ self.cpr_support = CPR_Support.NOT_SUPPORTED
+
+ if self.cpr_not_supported_callback:
+ # Make sure to call this callback in the main thread.
+ self.cpr_not_supported_callback()
+
+ get_app().create_background_task(timer())
+
+ def report_absolute_cursor_row(self, row: int) -> None:
+ """
+ To be called when we know the absolute cursor position.
+ (As an answer of a "Cursor Position Request" response.)
+ """
+ self.cpr_support = CPR_Support.SUPPORTED
+
+ # Calculate the amount of rows from the cursor position until the
+ # bottom of the terminal.
+ total_rows = self.output.get_size().rows
+ rows_below_cursor = total_rows - row + 1
+
+ # Set the minimum available height.
+ self._min_available_height = rows_below_cursor
+
+ # Pop and set waiting for CPR future.
+ try:
+ f = self._waiting_for_cpr_futures.popleft()
+ except IndexError:
+ pass # Received CPR response without having a CPR.
+ else:
+ f.set_result(None)
+
+ @property
+ def waiting_for_cpr(self) -> bool:
+ """
+ Waiting for CPR flag. True when we send the request, but didn't got a
+ response.
+ """
+ return bool(self._waiting_for_cpr_futures)
+
+ async def wait_for_cpr_responses(self, timeout: int = 1) -> None:
+ """
+ Wait for a CPR response.
+ """
+ cpr_futures = list(self._waiting_for_cpr_futures) # Make copy.
+
+ # When there are no CPRs in the queue. Don't do anything.
+ if not cpr_futures or self.cpr_support == CPR_Support.NOT_SUPPORTED:
+ return None
+
+ async def wait_for_responses() -> None:
+ for response_f in cpr_futures:
+ await response_f
+
+ async def wait_for_timeout() -> None:
+ await sleep(timeout)
+
+ # Got timeout, erase queue.
+ for response_f in cpr_futures:
+ response_f.cancel()
+ self._waiting_for_cpr_futures = deque()
+
+ tasks = {
+ ensure_future(wait_for_responses()),
+ ensure_future(wait_for_timeout()),
+ }
+ _, pending = await wait(tasks, return_when=FIRST_COMPLETED)
+ for task in pending:
+ task.cancel()
+
+ def render(
+ self, app: Application[Any], layout: Layout, is_done: bool = False
+ ) -> None:
+ """
+ Render the current interface to the output.
+
+ :param is_done: When True, put the cursor at the end of the interface. We
+ won't print any changes to this part.
+ """
+ output = self.output
+
+ # Enter alternate screen.
+ if self.full_screen and not self._in_alternate_screen:
+ self._in_alternate_screen = True
+ output.enter_alternate_screen()
+
+ # Enable bracketed paste.
+ if not self._bracketed_paste_enabled:
+ self.output.enable_bracketed_paste()
+ self._bracketed_paste_enabled = True
+
+ # Reset cursor key mode.
+ if not self._cursor_key_mode_reset:
+ self.output.reset_cursor_key_mode()
+ self._cursor_key_mode_reset = True
+
+ # Enable/disable mouse support.
+ needs_mouse_support = self.mouse_support()
+
+ if needs_mouse_support and not self._mouse_support_enabled:
+ output.enable_mouse_support()
+ self._mouse_support_enabled = True
+
+ elif not needs_mouse_support and self._mouse_support_enabled:
+ output.disable_mouse_support()
+ self._mouse_support_enabled = False
+
+ # Create screen and write layout to it.
+ size = output.get_size()
+ screen = Screen()
+ screen.show_cursor = False # Hide cursor by default, unless one of the
+ # containers decides to display it.
+ mouse_handlers = MouseHandlers()
+
+ # Calculate height.
+ if self.full_screen:
+ height = size.rows
+ elif is_done:
+ # When we are done, we don't necessary want to fill up until the bottom.
+ height = layout.container.preferred_height(
+ size.columns, size.rows
+ ).preferred
+ else:
+ last_height = self._last_screen.height if self._last_screen else 0
+ height = max(
+ self._min_available_height,
+ last_height,
+ layout.container.preferred_height(size.columns, size.rows).preferred,
+ )
+
+ height = min(height, size.rows)
+
+ # When the size changes, don't consider the previous screen.
+ if self._last_size != size:
+ self._last_screen = None
+
+ # When we render using another style or another color depth, do a full
+ # repaint. (Forget about the previous rendered screen.)
+ # (But note that we still use _last_screen to calculate the height.)
+ if (
+ self.style.invalidation_hash() != self._last_style_hash
+ or app.style_transformation.invalidation_hash()
+ != self._last_transformation_hash
+ or app.color_depth != self._last_color_depth
+ ):
+ self._last_screen = None
+ self._attrs_for_style = None
+ self._style_string_has_style = None
+
+ if self._attrs_for_style is None:
+ self._attrs_for_style = _StyleStringToAttrsCache(
+ self.style.get_attrs_for_style_str, app.style_transformation
+ )
+ if self._style_string_has_style is None:
+ self._style_string_has_style = _StyleStringHasStyleCache(
+ self._attrs_for_style
+ )
+
+ self._last_style_hash = self.style.invalidation_hash()
+ self._last_transformation_hash = app.style_transformation.invalidation_hash()
+ self._last_color_depth = app.color_depth
+
+ layout.container.write_to_screen(
+ screen,
+ mouse_handlers,
+ WritePosition(xpos=0, ypos=0, width=size.columns, height=height),
+ parent_style="",
+ erase_bg=False,
+ z_index=None,
+ )
+ screen.draw_all_floats()
+
+ # When grayed. Replace all styles in the new screen.
+ if app.exit_style:
+ screen.append_style_to_content(app.exit_style)
+
+ # Process diff and write to output.
+ self._cursor_pos, self._last_style = _output_screen_diff(
+ app,
+ output,
+ screen,
+ self._cursor_pos,
+ app.color_depth,
+ self._last_screen,
+ self._last_style,
+ is_done,
+ full_screen=self.full_screen,
+ attrs_for_style_string=self._attrs_for_style,
+ style_string_has_style=self._style_string_has_style,
+ size=size,
+ previous_width=(self._last_size.columns if self._last_size else 0),
+ )
+ self._last_screen = screen
+ self._last_size = size
+ self.mouse_handlers = mouse_handlers
+
+ # Handle cursor shapes.
+ new_cursor_shape = app.cursor.get_cursor_shape(app)
+ if (
+ self._last_cursor_shape is None
+ or self._last_cursor_shape != new_cursor_shape
+ ):
+ output.set_cursor_shape(new_cursor_shape)
+ self._last_cursor_shape = new_cursor_shape
+
+ # Flush buffered output.
+ output.flush()
+
+ # Set visible windows in layout.
+ app.layout.visible_windows = screen.visible_windows
+
+ if is_done:
+ self.reset()
+
+ def erase(self, leave_alternate_screen: bool = True) -> None:
+ """
+ Hide all output and put the cursor back at the first line. This is for
+ instance used for running a system command (while hiding the CLI) and
+ later resuming the same CLI.)
+
+ :param leave_alternate_screen: When True, and when inside an alternate
+ screen buffer, quit the alternate screen.
+ """
+ output = self.output
+
+ output.cursor_backward(self._cursor_pos.x)
+ output.cursor_up(self._cursor_pos.y)
+ output.erase_down()
+ output.reset_attributes()
+ output.enable_autowrap()
+
+ output.flush()
+
+ self.reset(leave_alternate_screen=leave_alternate_screen)
+
+ def clear(self) -> None:
+ """
+ Clear screen and go to 0,0
+ """
+ # Erase current output first.
+ self.erase()
+
+ # Send "Erase Screen" command and go to (0, 0).
+ output = self.output
+
+ output.erase_screen()
+ output.cursor_goto(0, 0)
+ output.flush()
+
+ self.request_absolute_cursor_position()
+
+
+def print_formatted_text(
+ output: Output,
+ formatted_text: AnyFormattedText,
+ style: BaseStyle,
+ style_transformation: StyleTransformation | None = None,
+ color_depth: ColorDepth | None = None,
+) -> None:
+ """
+ Print a list of (style_str, text) tuples in the given style to the output.
+ """
+ fragments = to_formatted_text(formatted_text)
+ style_transformation = style_transformation or DummyStyleTransformation()
+ color_depth = color_depth or output.get_default_color_depth()
+
+ # Reset first.
+ output.reset_attributes()
+ output.enable_autowrap()
+ last_attrs: Attrs | None = None
+
+ # Print all (style_str, text) tuples.
+ attrs_for_style_string = _StyleStringToAttrsCache(
+ style.get_attrs_for_style_str, style_transformation
+ )
+
+ for style_str, text, *_ in fragments:
+ attrs = attrs_for_style_string[style_str]
+
+ # Set style attributes if something changed.
+ if attrs != last_attrs:
+ if attrs:
+ output.set_attributes(attrs, color_depth)
+ else:
+ output.reset_attributes()
+ last_attrs = attrs
+
+ # Print escape sequences as raw output
+ if "[ZeroWidthEscape]" in style_str:
+ output.write_raw(text)
+ else:
+ # Eliminate carriage returns
+ text = text.replace("\r", "")
+ # Insert a carriage return before every newline (important when the
+ # front-end is a telnet client).
+ text = text.replace("\n", "\r\n")
+ output.write(text)
+
+ # Reset again.
+ output.reset_attributes()
+ output.flush()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/search.py b/lib/python3.12/site-packages/prompt_toolkit/search.py
new file mode 100644
index 0000000000000000000000000000000000000000..d1cf7acbac6c25045c7cb46e29d41a20a4b2be03
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/search.py
@@ -0,0 +1,226 @@
+"""
+Search operations.
+
+For the key bindings implementation with attached filters, check
+`prompt_toolkit.key_binding.bindings.search`. (Use these for new key bindings
+instead of calling these function directly.)
+"""
+
+from __future__ import annotations
+
+from enum import Enum
+from typing import TYPE_CHECKING
+
+from .application.current import get_app
+from .filters import FilterOrBool, is_searching, to_filter
+from .key_binding.vi_state import InputMode
+
+if TYPE_CHECKING:
+ from prompt_toolkit.layout.controls import BufferControl, SearchBufferControl
+ from prompt_toolkit.layout.layout import Layout
+
+__all__ = [
+ "SearchDirection",
+ "start_search",
+ "stop_search",
+]
+
+
+class SearchDirection(Enum):
+ FORWARD = "FORWARD"
+ BACKWARD = "BACKWARD"
+
+
+class SearchState:
+ """
+ A search 'query', associated with a search field (like a SearchToolbar).
+
+ Every searchable `BufferControl` points to a `search_buffer_control`
+ (another `BufferControls`) which represents the search field. The
+ `SearchState` attached to that search field is used for storing the current
+ search query.
+
+ It is possible to have one searchfield for multiple `BufferControls`. In
+ that case, they'll share the same `SearchState`.
+ If there are multiple `BufferControls` that display the same `Buffer`, then
+ they can have a different `SearchState` each (if they have a different
+ search control).
+ """
+
+ __slots__ = ("text", "direction", "ignore_case")
+
+ def __init__(
+ self,
+ text: str = "",
+ direction: SearchDirection = SearchDirection.FORWARD,
+ ignore_case: FilterOrBool = False,
+ ) -> None:
+ self.text = text
+ self.direction = direction
+ self.ignore_case = to_filter(ignore_case)
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}({self.text!r}, direction={self.direction!r}, ignore_case={self.ignore_case!r})"
+
+ def __invert__(self) -> SearchState:
+ """
+ Create a new SearchState where backwards becomes forwards and the other
+ way around.
+ """
+ if self.direction == SearchDirection.BACKWARD:
+ direction = SearchDirection.FORWARD
+ else:
+ direction = SearchDirection.BACKWARD
+
+ return SearchState(
+ text=self.text, direction=direction, ignore_case=self.ignore_case
+ )
+
+
+def start_search(
+ buffer_control: BufferControl | None = None,
+ direction: SearchDirection = SearchDirection.FORWARD,
+) -> None:
+ """
+ Start search through the given `buffer_control` using the
+ `search_buffer_control`.
+
+ :param buffer_control: Start search for this `BufferControl`. If not given,
+ search through the current control.
+ """
+ from prompt_toolkit.layout.controls import BufferControl
+
+ assert buffer_control is None or isinstance(buffer_control, BufferControl)
+
+ layout = get_app().layout
+
+ # When no control is given, use the current control if that's a BufferControl.
+ if buffer_control is None:
+ if not isinstance(layout.current_control, BufferControl):
+ return
+ buffer_control = layout.current_control
+
+ # Only if this control is searchable.
+ search_buffer_control = buffer_control.search_buffer_control
+
+ if search_buffer_control:
+ buffer_control.search_state.direction = direction
+
+ # Make sure to focus the search BufferControl
+ layout.focus(search_buffer_control)
+
+ # Remember search link.
+ layout.search_links[search_buffer_control] = buffer_control
+
+ # If we're in Vi mode, make sure to go into insert mode.
+ get_app().vi_state.input_mode = InputMode.INSERT
+
+
+def stop_search(buffer_control: BufferControl | None = None) -> None:
+ """
+ Stop search through the given `buffer_control`.
+ """
+ layout = get_app().layout
+
+ if buffer_control is None:
+ buffer_control = layout.search_target_buffer_control
+ if buffer_control is None:
+ # (Should not happen, but possible when `stop_search` is called
+ # when we're not searching.)
+ return
+ search_buffer_control = buffer_control.search_buffer_control
+ else:
+ assert buffer_control in layout.search_links.values()
+ search_buffer_control = _get_reverse_search_links(layout)[buffer_control]
+
+ # Focus the original buffer again.
+ layout.focus(buffer_control)
+
+ if search_buffer_control is not None:
+ # Remove the search link.
+ del layout.search_links[search_buffer_control]
+
+ # Reset content of search control.
+ search_buffer_control.buffer.reset()
+
+ # If we're in Vi mode, go back to navigation mode.
+ get_app().vi_state.input_mode = InputMode.NAVIGATION
+
+
+def do_incremental_search(direction: SearchDirection, count: int = 1) -> None:
+ """
+ Apply search, but keep search buffer focused.
+ """
+ assert is_searching()
+
+ layout = get_app().layout
+
+ # Only search if the current control is a `BufferControl`.
+ from prompt_toolkit.layout.controls import BufferControl
+
+ search_control = layout.current_control
+ if not isinstance(search_control, BufferControl):
+ return
+
+ prev_control = layout.search_target_buffer_control
+ if prev_control is None:
+ return
+ search_state = prev_control.search_state
+
+ # Update search_state.
+ direction_changed = search_state.direction != direction
+
+ search_state.text = search_control.buffer.text
+ search_state.direction = direction
+
+ # Apply search to current buffer.
+ if not direction_changed:
+ prev_control.buffer.apply_search(
+ search_state, include_current_position=False, count=count
+ )
+
+
+def accept_search() -> None:
+ """
+ Accept current search query. Focus original `BufferControl` again.
+ """
+ layout = get_app().layout
+
+ search_control = layout.current_control
+ target_buffer_control = layout.search_target_buffer_control
+
+ from prompt_toolkit.layout.controls import BufferControl
+
+ if not isinstance(search_control, BufferControl):
+ return
+ if target_buffer_control is None:
+ return
+
+ search_state = target_buffer_control.search_state
+
+ # Update search state.
+ if search_control.buffer.text:
+ search_state.text = search_control.buffer.text
+
+ # Apply search.
+ target_buffer_control.buffer.apply_search(
+ search_state, include_current_position=True
+ )
+
+ # Add query to history of search line.
+ search_control.buffer.append_to_history()
+
+ # Stop search and focus previous control again.
+ stop_search(target_buffer_control)
+
+
+def _get_reverse_search_links(
+ layout: Layout,
+) -> dict[BufferControl, SearchBufferControl]:
+ """
+ Return mapping from BufferControl to SearchBufferControl.
+ """
+ return {
+ buffer_control: search_buffer_control
+ for search_buffer_control, buffer_control in layout.search_links.items()
+ }
diff --git a/lib/python3.12/site-packages/prompt_toolkit/selection.py b/lib/python3.12/site-packages/prompt_toolkit/selection.py
new file mode 100644
index 0000000000000000000000000000000000000000..ff88535b0f71b6f5da05f98e9174e406a7e89dac
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/selection.py
@@ -0,0 +1,58 @@
+"""
+Data structures for the selection.
+"""
+
+from __future__ import annotations
+
+from enum import Enum
+
+__all__ = [
+ "SelectionType",
+ "PasteMode",
+ "SelectionState",
+]
+
+
+class SelectionType(Enum):
+ """
+ Type of selection.
+ """
+
+ #: Characters. (Visual in Vi.)
+ CHARACTERS = "CHARACTERS"
+
+ #: Whole lines. (Visual-Line in Vi.)
+ LINES = "LINES"
+
+ #: A block selection. (Visual-Block in Vi.)
+ BLOCK = "BLOCK"
+
+
+class PasteMode(Enum):
+ EMACS = "EMACS" # Yank like emacs.
+ VI_AFTER = "VI_AFTER" # When pressing 'p' in Vi.
+ VI_BEFORE = "VI_BEFORE" # When pressing 'P' in Vi.
+
+
+class SelectionState:
+ """
+ State of the current selection.
+
+ :param original_cursor_position: int
+ :param type: :class:`~.SelectionType`
+ """
+
+ def __init__(
+ self,
+ original_cursor_position: int = 0,
+ type: SelectionType = SelectionType.CHARACTERS,
+ ) -> None:
+ self.original_cursor_position = original_cursor_position
+ self.type = type
+ self.shift_mode = False
+
+ def enter_shift_mode(self) -> None:
+ self.shift_mode = True
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(original_cursor_position={self.original_cursor_position!r}, type={self.type!r})"
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__init__.py b/lib/python3.12/site-packages/prompt_toolkit/styles/__init__.py
new file mode 100644
index 0000000000000000000000000000000000000000..39e97cac8ef0099f7beb9f408c2404ed8da4a8f7
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/__init__.py
@@ -0,0 +1,67 @@
+"""
+Styling for prompt_toolkit applications.
+"""
+
+from __future__ import annotations
+
+from .base import (
+ ANSI_COLOR_NAMES,
+ DEFAULT_ATTRS,
+ Attrs,
+ BaseStyle,
+ DummyStyle,
+ DynamicStyle,
+)
+from .defaults import default_pygments_style, default_ui_style
+from .named_colors import NAMED_COLORS
+from .pygments import (
+ pygments_token_to_classname,
+ style_from_pygments_cls,
+ style_from_pygments_dict,
+)
+from .style import Priority, Style, merge_styles, parse_color
+from .style_transformation import (
+ AdjustBrightnessStyleTransformation,
+ ConditionalStyleTransformation,
+ DummyStyleTransformation,
+ DynamicStyleTransformation,
+ ReverseStyleTransformation,
+ SetDefaultColorStyleTransformation,
+ StyleTransformation,
+ SwapLightAndDarkStyleTransformation,
+ merge_style_transformations,
+)
+
+__all__ = [
+ # Base.
+ "Attrs",
+ "DEFAULT_ATTRS",
+ "ANSI_COLOR_NAMES",
+ "BaseStyle",
+ "DummyStyle",
+ "DynamicStyle",
+ # Defaults.
+ "default_ui_style",
+ "default_pygments_style",
+ # Style.
+ "Style",
+ "Priority",
+ "merge_styles",
+ "parse_color",
+ # Style transformation.
+ "StyleTransformation",
+ "SwapLightAndDarkStyleTransformation",
+ "ReverseStyleTransformation",
+ "SetDefaultColorStyleTransformation",
+ "AdjustBrightnessStyleTransformation",
+ "DummyStyleTransformation",
+ "ConditionalStyleTransformation",
+ "DynamicStyleTransformation",
+ "merge_style_transformations",
+ # Pygments.
+ "style_from_pygments_cls",
+ "style_from_pygments_dict",
+ "pygments_token_to_classname",
+ # Named colors.
+ "NAMED_COLORS",
+]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..61c1668c7c6544e4c32f9ba5bbc4f729eaf64328
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/base.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/base.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..a818e9378d83b657bb3027c75f6b1f14e9f2c157
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/base.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/defaults.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/defaults.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..9375038213da0ff95ea67150c5b1e6814a1e1641
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/defaults.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/named_colors.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/named_colors.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..afba533d06547d5e8d6a1fc951b7e636645a6684
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/named_colors.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/pygments.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/pygments.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..6cfa9c22beb54e7f35a0c2cd7dcd0acf06ddad54
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/pygments.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/style.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/style.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..03ae5c5b635d437ded82211e6e62e9e3737de5c2
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/style.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/style_transformation.cpython-312.pyc b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/style_transformation.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..861e897ab1bc59c93f585a0d56a0f8a00cbf784c
Binary files /dev/null and b/lib/python3.12/site-packages/prompt_toolkit/styles/__pycache__/style_transformation.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/base.py b/lib/python3.12/site-packages/prompt_toolkit/styles/base.py
new file mode 100644
index 0000000000000000000000000000000000000000..1e044679f0f0c19bb4aa8c1c605ebe8f58d6b95c
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/base.py
@@ -0,0 +1,188 @@
+"""
+The base classes for the styling.
+"""
+
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from typing import Callable, Hashable, NamedTuple
+
+__all__ = [
+ "Attrs",
+ "DEFAULT_ATTRS",
+ "ANSI_COLOR_NAMES",
+ "ANSI_COLOR_NAMES_ALIASES",
+ "BaseStyle",
+ "DummyStyle",
+ "DynamicStyle",
+]
+
+
+#: Style attributes.
+class Attrs(NamedTuple):
+ color: str | None
+ bgcolor: str | None
+ bold: bool | None
+ underline: bool | None
+ strike: bool | None
+ italic: bool | None
+ blink: bool | None
+ reverse: bool | None
+ hidden: bool | None
+ dim: bool | None
+
+
+"""
+:param color: Hexadecimal string. E.g. '000000' or Ansi color name: e.g. 'ansiblue'
+:param bgcolor: Hexadecimal string. E.g. 'ffffff' or Ansi color name: e.g. 'ansired'
+:param bold: Boolean
+:param underline: Boolean
+:param strike: Boolean
+:param italic: Boolean
+:param blink: Boolean
+:param reverse: Boolean
+:param hidden: Boolean
+:param dim: Boolean
+"""
+
+#: The default `Attrs`.
+DEFAULT_ATTRS = Attrs(
+ color="",
+ bgcolor="",
+ bold=False,
+ underline=False,
+ strike=False,
+ italic=False,
+ blink=False,
+ reverse=False,
+ hidden=False,
+ dim=False,
+)
+
+
+#: ``Attrs.bgcolor/fgcolor`` can be in either 'ffffff' format, or can be any of
+#: the following in case we want to take colors from the 8/16 color palette.
+#: Usually, in that case, the terminal application allows to configure the RGB
+#: values for these names.
+#: ISO 6429 colors
+ANSI_COLOR_NAMES = [
+ "ansidefault",
+ # Low intensity, dark. (One or two components 0x80, the other 0x00.)
+ "ansiblack",
+ "ansired",
+ "ansigreen",
+ "ansiyellow",
+ "ansiblue",
+ "ansimagenta",
+ "ansicyan",
+ "ansigray",
+ # High intensity, bright. (One or two components 0xff, the other 0x00. Not supported everywhere.)
+ "ansibrightblack",
+ "ansibrightred",
+ "ansibrightgreen",
+ "ansibrightyellow",
+ "ansibrightblue",
+ "ansibrightmagenta",
+ "ansibrightcyan",
+ "ansiwhite",
+]
+
+
+# People don't use the same ANSI color names everywhere. In prompt_toolkit 1.0
+# we used some unconventional names (which were contributed like that to
+# Pygments). This is fixed now, but we still support the old names.
+
+# The table below maps the old aliases to the current names.
+ANSI_COLOR_NAMES_ALIASES: dict[str, str] = {
+ "ansidarkgray": "ansibrightblack",
+ "ansiteal": "ansicyan",
+ "ansiturquoise": "ansibrightcyan",
+ "ansibrown": "ansiyellow",
+ "ansipurple": "ansimagenta",
+ "ansifuchsia": "ansibrightmagenta",
+ "ansilightgray": "ansigray",
+ "ansidarkred": "ansired",
+ "ansidarkgreen": "ansigreen",
+ "ansidarkblue": "ansiblue",
+}
+assert set(ANSI_COLOR_NAMES_ALIASES.values()).issubset(set(ANSI_COLOR_NAMES))
+assert not (set(ANSI_COLOR_NAMES_ALIASES.keys()) & set(ANSI_COLOR_NAMES))
+
+
+class BaseStyle(metaclass=ABCMeta):
+ """
+ Abstract base class for prompt_toolkit styles.
+ """
+
+ @abstractmethod
+ def get_attrs_for_style_str(
+ self, style_str: str, default: Attrs = DEFAULT_ATTRS
+ ) -> Attrs:
+ """
+ Return :class:`.Attrs` for the given style string.
+
+ :param style_str: The style string. This can contain inline styling as
+ well as classnames (e.g. "class:title").
+ :param default: `Attrs` to be used if no styling was defined.
+ """
+
+ @property
+ @abstractmethod
+ def style_rules(self) -> list[tuple[str, str]]:
+ """
+ The list of style rules, used to create this style.
+ (Required for `DynamicStyle` and `_MergedStyle` to work.)
+ """
+ return []
+
+ @abstractmethod
+ def invalidation_hash(self) -> Hashable:
+ """
+ Invalidation hash for the style. When this changes over time, the
+ renderer knows that something in the style changed, and that everything
+ has to be redrawn.
+ """
+
+
+class DummyStyle(BaseStyle):
+ """
+ A style that doesn't style anything.
+ """
+
+ def get_attrs_for_style_str(
+ self, style_str: str, default: Attrs = DEFAULT_ATTRS
+ ) -> Attrs:
+ return default
+
+ def invalidation_hash(self) -> Hashable:
+ return 1 # Always the same value.
+
+ @property
+ def style_rules(self) -> list[tuple[str, str]]:
+ return []
+
+
+class DynamicStyle(BaseStyle):
+ """
+ Style class that can dynamically returns an other Style.
+
+ :param get_style: Callable that returns a :class:`.Style` instance.
+ """
+
+ def __init__(self, get_style: Callable[[], BaseStyle | None]):
+ self.get_style = get_style
+ self._dummy = DummyStyle()
+
+ def get_attrs_for_style_str(
+ self, style_str: str, default: Attrs = DEFAULT_ATTRS
+ ) -> Attrs:
+ style = self.get_style() or self._dummy
+
+ return style.get_attrs_for_style_str(style_str, default)
+
+ def invalidation_hash(self) -> Hashable:
+ return (self.get_style() or self._dummy).invalidation_hash()
+
+ @property
+ def style_rules(self) -> list[tuple[str, str]]:
+ return (self.get_style() or self._dummy).style_rules
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/defaults.py b/lib/python3.12/site-packages/prompt_toolkit/styles/defaults.py
new file mode 100644
index 0000000000000000000000000000000000000000..2faba94378ef30cb47781ccaa6176992361b8d58
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/defaults.py
@@ -0,0 +1,236 @@
+"""
+The default styling.
+"""
+
+from __future__ import annotations
+
+from prompt_toolkit.cache import memoized
+
+from .base import ANSI_COLOR_NAMES, BaseStyle
+from .named_colors import NAMED_COLORS
+from .style import Style, merge_styles
+
+__all__ = [
+ "default_ui_style",
+ "default_pygments_style",
+]
+
+#: Default styling. Mapping from classnames to their style definition.
+PROMPT_TOOLKIT_STYLE = [
+ # Highlighting of search matches in document.
+ ("search", "bg:ansibrightyellow ansiblack"),
+ ("search.current", ""),
+ # Incremental search.
+ ("incsearch", ""),
+ ("incsearch.current", "reverse"),
+ # Highlighting of select text in document.
+ ("selected", "reverse"),
+ ("cursor-column", "bg:#dddddd"),
+ ("cursor-line", "underline"),
+ ("color-column", "bg:#ccaacc"),
+ # Highlighting of matching brackets.
+ ("matching-bracket", ""),
+ ("matching-bracket.other", "#000000 bg:#aacccc"),
+ ("matching-bracket.cursor", "#ff8888 bg:#880000"),
+ # Styling of other cursors, in case of block editing.
+ ("multiple-cursors", "#000000 bg:#ccccaa"),
+ # Line numbers.
+ ("line-number", "#888888"),
+ ("line-number.current", "bold"),
+ ("tilde", "#8888ff"),
+ # Default prompt.
+ ("prompt", ""),
+ ("prompt.arg", "noinherit"),
+ ("prompt.arg.text", ""),
+ ("prompt.search", "noinherit"),
+ ("prompt.search.text", ""),
+ # Search toolbar.
+ ("search-toolbar", "bold"),
+ ("search-toolbar.text", "nobold"),
+ # System toolbar
+ ("system-toolbar", "bold"),
+ ("system-toolbar.text", "nobold"),
+ # "arg" toolbar.
+ ("arg-toolbar", "bold"),
+ ("arg-toolbar.text", "nobold"),
+ # Validation toolbar.
+ ("validation-toolbar", "bg:#550000 #ffffff"),
+ ("window-too-small", "bg:#550000 #ffffff"),
+ # Completions toolbar.
+ ("completion-toolbar", "bg:#bbbbbb #000000"),
+ ("completion-toolbar.arrow", "bg:#bbbbbb #000000 bold"),
+ ("completion-toolbar.completion", "bg:#bbbbbb #000000"),
+ ("completion-toolbar.completion.current", "bg:#444444 #ffffff"),
+ # Completions menu.
+ ("completion-menu", "bg:#bbbbbb #000000"),
+ ("completion-menu.completion", ""),
+ # (Note: for the current completion, we use 'reverse' on top of fg/bg
+ # colors. This is to have proper rendering with NO_COLOR=1).
+ ("completion-menu.completion.current", "fg:#888888 bg:#ffffff reverse"),
+ ("completion-menu.meta.completion", "bg:#999999 #000000"),
+ ("completion-menu.meta.completion.current", "bg:#aaaaaa #000000"),
+ ("completion-menu.multi-column-meta", "bg:#aaaaaa #000000"),
+ # Fuzzy matches in completion menu (for FuzzyCompleter).
+ ("completion-menu.completion fuzzymatch.outside", "fg:#444444"),
+ ("completion-menu.completion fuzzymatch.inside", "bold"),
+ ("completion-menu.completion fuzzymatch.inside.character", "underline"),
+ ("completion-menu.completion.current fuzzymatch.outside", "fg:default"),
+ ("completion-menu.completion.current fuzzymatch.inside", "nobold"),
+ # Styling of readline-like completions.
+ ("readline-like-completions", ""),
+ ("readline-like-completions.completion", ""),
+ ("readline-like-completions.completion fuzzymatch.outside", "#888888"),
+ ("readline-like-completions.completion fuzzymatch.inside", ""),
+ ("readline-like-completions.completion fuzzymatch.inside.character", "underline"),
+ # Scrollbars.
+ ("scrollbar.background", "bg:#aaaaaa"),
+ ("scrollbar.button", "bg:#444444"),
+ ("scrollbar.arrow", "noinherit bold"),
+ # Start/end of scrollbars. Adding 'underline' here provides a nice little
+ # detail to the progress bar, but it doesn't look good on all terminals.
+ # ('scrollbar.start', 'underline #ffffff'),
+ # ('scrollbar.end', 'underline #000000'),
+ # Auto suggestion text.
+ ("auto-suggestion", "#666666"),
+ # Trailing whitespace and tabs.
+ ("trailing-whitespace", "#999999"),
+ ("tab", "#999999"),
+ # When Control-C/D has been pressed. Grayed.
+ ("aborting", "#888888 bg:default noreverse noitalic nounderline noblink"),
+ ("exiting", "#888888 bg:default noreverse noitalic nounderline noblink"),
+ # Entering a Vi digraph.
+ ("digraph", "#4444ff"),
+ # Control characters, like ^C, ^X.
+ ("control-character", "ansiblue"),
+ # Non-breaking space.
+ ("nbsp", "underline ansiyellow"),
+ # Default styling of HTML elements.
+ ("i", "italic"),
+ ("u", "underline"),
+ ("s", "strike"),
+ ("b", "bold"),
+ ("em", "italic"),
+ ("strong", "bold"),
+ ("del", "strike"),
+ ("hidden", "hidden"),
+ # It should be possible to use the style names in HTML.
+ # ... or ....
+ ("italic", "italic"),
+ ("underline", "underline"),
+ ("strike", "strike"),
+ ("bold", "bold"),
+ ("reverse", "reverse"),
+ ("noitalic", "noitalic"),
+ ("nounderline", "nounderline"),
+ ("nostrike", "nostrike"),
+ ("nobold", "nobold"),
+ ("noreverse", "noreverse"),
+ # Prompt bottom toolbar
+ ("bottom-toolbar", "reverse"),
+]
+
+
+# Style that will turn for instance the class 'red' into 'red'.
+COLORS_STYLE = [(name, "fg:" + name) for name in ANSI_COLOR_NAMES] + [
+ (name.lower(), "fg:" + name) for name in NAMED_COLORS
+]
+
+
+WIDGETS_STYLE = [
+ # Dialog windows.
+ ("dialog", "bg:#4444ff"),
+ ("dialog.body", "bg:#ffffff #000000"),
+ ("dialog.body text-area", "bg:#cccccc"),
+ ("dialog.body text-area last-line", "underline"),
+ ("dialog frame.label", "#ff0000 bold"),
+ # Scrollbars in dialogs.
+ ("dialog.body scrollbar.background", ""),
+ ("dialog.body scrollbar.button", "bg:#000000"),
+ ("dialog.body scrollbar.arrow", ""),
+ ("dialog.body scrollbar.start", "nounderline"),
+ ("dialog.body scrollbar.end", "nounderline"),
+ # Buttons.
+ ("button", ""),
+ ("button.arrow", "bold"),
+ ("button.focused", "bg:#aa0000 #ffffff"),
+ # Menu bars.
+ ("menu-bar", "bg:#aaaaaa #000000"),
+ ("menu-bar.selected-item", "bg:#ffffff #000000"),
+ ("menu", "bg:#888888 #ffffff"),
+ ("menu.border", "#aaaaaa"),
+ ("menu.border shadow", "#444444"),
+ # Shadows.
+ ("dialog shadow", "bg:#000088"),
+ ("dialog.body shadow", "bg:#aaaaaa"),
+ ("progress-bar", "bg:#000088"),
+ ("progress-bar.used", "bg:#ff0000"),
+]
+
+
+# The default Pygments style, include this by default in case a Pygments lexer
+# is used.
+PYGMENTS_DEFAULT_STYLE = {
+ "pygments.whitespace": "#bbbbbb",
+ "pygments.comment": "italic #408080",
+ "pygments.comment.preproc": "noitalic #bc7a00",
+ "pygments.keyword": "bold #008000",
+ "pygments.keyword.pseudo": "nobold",
+ "pygments.keyword.type": "nobold #b00040",
+ "pygments.operator": "#666666",
+ "pygments.operator.word": "bold #aa22ff",
+ "pygments.name.builtin": "#008000",
+ "pygments.name.function": "#0000ff",
+ "pygments.name.class": "bold #0000ff",
+ "pygments.name.namespace": "bold #0000ff",
+ "pygments.name.exception": "bold #d2413a",
+ "pygments.name.variable": "#19177c",
+ "pygments.name.constant": "#880000",
+ "pygments.name.label": "#a0a000",
+ "pygments.name.entity": "bold #999999",
+ "pygments.name.attribute": "#7d9029",
+ "pygments.name.tag": "bold #008000",
+ "pygments.name.decorator": "#aa22ff",
+ # Note: In Pygments, Token.String is an alias for Token.Literal.String,
+ # and Token.Number as an alias for Token.Literal.Number.
+ "pygments.literal.string": "#ba2121",
+ "pygments.literal.string.doc": "italic",
+ "pygments.literal.string.interpol": "bold #bb6688",
+ "pygments.literal.string.escape": "bold #bb6622",
+ "pygments.literal.string.regex": "#bb6688",
+ "pygments.literal.string.symbol": "#19177c",
+ "pygments.literal.string.other": "#008000",
+ "pygments.literal.number": "#666666",
+ "pygments.generic.heading": "bold #000080",
+ "pygments.generic.subheading": "bold #800080",
+ "pygments.generic.deleted": "#a00000",
+ "pygments.generic.inserted": "#00a000",
+ "pygments.generic.error": "#ff0000",
+ "pygments.generic.emph": "italic",
+ "pygments.generic.strong": "bold",
+ "pygments.generic.prompt": "bold #000080",
+ "pygments.generic.output": "#888",
+ "pygments.generic.traceback": "#04d",
+ "pygments.error": "border:#ff0000",
+}
+
+
+@memoized()
+def default_ui_style() -> BaseStyle:
+ """
+ Create a default `Style` object.
+ """
+ return merge_styles(
+ [
+ Style(PROMPT_TOOLKIT_STYLE),
+ Style(COLORS_STYLE),
+ Style(WIDGETS_STYLE),
+ ]
+ )
+
+
+@memoized()
+def default_pygments_style() -> Style:
+ """
+ Create a `Style` object that contains the default Pygments style.
+ """
+ return Style.from_dict(PYGMENTS_DEFAULT_STYLE)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/named_colors.py b/lib/python3.12/site-packages/prompt_toolkit/styles/named_colors.py
new file mode 100644
index 0000000000000000000000000000000000000000..b6290a4e8fe6f99fbb91770f6fc68f5a51b4fd3b
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/named_colors.py
@@ -0,0 +1,162 @@
+"""
+All modern web browsers support these 140 color names.
+Taken from: https://www.w3schools.com/colors/colors_names.asp
+"""
+
+from __future__ import annotations
+
+__all__ = [
+ "NAMED_COLORS",
+]
+
+
+NAMED_COLORS: dict[str, str] = {
+ "AliceBlue": "#f0f8ff",
+ "AntiqueWhite": "#faebd7",
+ "Aqua": "#00ffff",
+ "Aquamarine": "#7fffd4",
+ "Azure": "#f0ffff",
+ "Beige": "#f5f5dc",
+ "Bisque": "#ffe4c4",
+ "Black": "#000000",
+ "BlanchedAlmond": "#ffebcd",
+ "Blue": "#0000ff",
+ "BlueViolet": "#8a2be2",
+ "Brown": "#a52a2a",
+ "BurlyWood": "#deb887",
+ "CadetBlue": "#5f9ea0",
+ "Chartreuse": "#7fff00",
+ "Chocolate": "#d2691e",
+ "Coral": "#ff7f50",
+ "CornflowerBlue": "#6495ed",
+ "Cornsilk": "#fff8dc",
+ "Crimson": "#dc143c",
+ "Cyan": "#00ffff",
+ "DarkBlue": "#00008b",
+ "DarkCyan": "#008b8b",
+ "DarkGoldenRod": "#b8860b",
+ "DarkGray": "#a9a9a9",
+ "DarkGreen": "#006400",
+ "DarkGrey": "#a9a9a9",
+ "DarkKhaki": "#bdb76b",
+ "DarkMagenta": "#8b008b",
+ "DarkOliveGreen": "#556b2f",
+ "DarkOrange": "#ff8c00",
+ "DarkOrchid": "#9932cc",
+ "DarkRed": "#8b0000",
+ "DarkSalmon": "#e9967a",
+ "DarkSeaGreen": "#8fbc8f",
+ "DarkSlateBlue": "#483d8b",
+ "DarkSlateGray": "#2f4f4f",
+ "DarkSlateGrey": "#2f4f4f",
+ "DarkTurquoise": "#00ced1",
+ "DarkViolet": "#9400d3",
+ "DeepPink": "#ff1493",
+ "DeepSkyBlue": "#00bfff",
+ "DimGray": "#696969",
+ "DimGrey": "#696969",
+ "DodgerBlue": "#1e90ff",
+ "FireBrick": "#b22222",
+ "FloralWhite": "#fffaf0",
+ "ForestGreen": "#228b22",
+ "Fuchsia": "#ff00ff",
+ "Gainsboro": "#dcdcdc",
+ "GhostWhite": "#f8f8ff",
+ "Gold": "#ffd700",
+ "GoldenRod": "#daa520",
+ "Gray": "#808080",
+ "Green": "#008000",
+ "GreenYellow": "#adff2f",
+ "Grey": "#808080",
+ "HoneyDew": "#f0fff0",
+ "HotPink": "#ff69b4",
+ "IndianRed": "#cd5c5c",
+ "Indigo": "#4b0082",
+ "Ivory": "#fffff0",
+ "Khaki": "#f0e68c",
+ "Lavender": "#e6e6fa",
+ "LavenderBlush": "#fff0f5",
+ "LawnGreen": "#7cfc00",
+ "LemonChiffon": "#fffacd",
+ "LightBlue": "#add8e6",
+ "LightCoral": "#f08080",
+ "LightCyan": "#e0ffff",
+ "LightGoldenRodYellow": "#fafad2",
+ "LightGray": "#d3d3d3",
+ "LightGreen": "#90ee90",
+ "LightGrey": "#d3d3d3",
+ "LightPink": "#ffb6c1",
+ "LightSalmon": "#ffa07a",
+ "LightSeaGreen": "#20b2aa",
+ "LightSkyBlue": "#87cefa",
+ "LightSlateGray": "#778899",
+ "LightSlateGrey": "#778899",
+ "LightSteelBlue": "#b0c4de",
+ "LightYellow": "#ffffe0",
+ "Lime": "#00ff00",
+ "LimeGreen": "#32cd32",
+ "Linen": "#faf0e6",
+ "Magenta": "#ff00ff",
+ "Maroon": "#800000",
+ "MediumAquaMarine": "#66cdaa",
+ "MediumBlue": "#0000cd",
+ "MediumOrchid": "#ba55d3",
+ "MediumPurple": "#9370db",
+ "MediumSeaGreen": "#3cb371",
+ "MediumSlateBlue": "#7b68ee",
+ "MediumSpringGreen": "#00fa9a",
+ "MediumTurquoise": "#48d1cc",
+ "MediumVioletRed": "#c71585",
+ "MidnightBlue": "#191970",
+ "MintCream": "#f5fffa",
+ "MistyRose": "#ffe4e1",
+ "Moccasin": "#ffe4b5",
+ "NavajoWhite": "#ffdead",
+ "Navy": "#000080",
+ "OldLace": "#fdf5e6",
+ "Olive": "#808000",
+ "OliveDrab": "#6b8e23",
+ "Orange": "#ffa500",
+ "OrangeRed": "#ff4500",
+ "Orchid": "#da70d6",
+ "PaleGoldenRod": "#eee8aa",
+ "PaleGreen": "#98fb98",
+ "PaleTurquoise": "#afeeee",
+ "PaleVioletRed": "#db7093",
+ "PapayaWhip": "#ffefd5",
+ "PeachPuff": "#ffdab9",
+ "Peru": "#cd853f",
+ "Pink": "#ffc0cb",
+ "Plum": "#dda0dd",
+ "PowderBlue": "#b0e0e6",
+ "Purple": "#800080",
+ "RebeccaPurple": "#663399",
+ "Red": "#ff0000",
+ "RosyBrown": "#bc8f8f",
+ "RoyalBlue": "#4169e1",
+ "SaddleBrown": "#8b4513",
+ "Salmon": "#fa8072",
+ "SandyBrown": "#f4a460",
+ "SeaGreen": "#2e8b57",
+ "SeaShell": "#fff5ee",
+ "Sienna": "#a0522d",
+ "Silver": "#c0c0c0",
+ "SkyBlue": "#87ceeb",
+ "SlateBlue": "#6a5acd",
+ "SlateGray": "#708090",
+ "SlateGrey": "#708090",
+ "Snow": "#fffafa",
+ "SpringGreen": "#00ff7f",
+ "SteelBlue": "#4682b4",
+ "Tan": "#d2b48c",
+ "Teal": "#008080",
+ "Thistle": "#d8bfd8",
+ "Tomato": "#ff6347",
+ "Turquoise": "#40e0d0",
+ "Violet": "#ee82ee",
+ "Wheat": "#f5deb3",
+ "White": "#ffffff",
+ "WhiteSmoke": "#f5f5f5",
+ "Yellow": "#ffff00",
+ "YellowGreen": "#9acd32",
+}
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/pygments.py b/lib/python3.12/site-packages/prompt_toolkit/styles/pygments.py
new file mode 100644
index 0000000000000000000000000000000000000000..c0f603108fb63817baef214ce413089d7cdda470
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/pygments.py
@@ -0,0 +1,70 @@
+"""
+Adaptor for building prompt_toolkit styles, starting from a Pygments style.
+
+Usage::
+
+ from pygments.styles.tango import TangoStyle
+ style = style_from_pygments_cls(pygments_style_cls=TangoStyle)
+"""
+
+from __future__ import annotations
+
+from typing import TYPE_CHECKING
+
+from .style import Style
+
+if TYPE_CHECKING:
+ from pygments.style import Style as PygmentsStyle
+ from pygments.token import Token
+
+
+__all__ = [
+ "style_from_pygments_cls",
+ "style_from_pygments_dict",
+ "pygments_token_to_classname",
+]
+
+
+def style_from_pygments_cls(pygments_style_cls: type[PygmentsStyle]) -> Style:
+ """
+ Shortcut to create a :class:`.Style` instance from a Pygments style class
+ and a style dictionary.
+
+ Example::
+
+ from prompt_toolkit.styles.from_pygments import style_from_pygments_cls
+ from pygments.styles import get_style_by_name
+ style = style_from_pygments_cls(get_style_by_name('monokai'))
+
+ :param pygments_style_cls: Pygments style class to start from.
+ """
+ # Import inline.
+ from pygments.style import Style as PygmentsStyle
+
+ assert issubclass(pygments_style_cls, PygmentsStyle)
+
+ return style_from_pygments_dict(pygments_style_cls.styles)
+
+
+def style_from_pygments_dict(pygments_dict: dict[Token, str]) -> Style:
+ """
+ Create a :class:`.Style` instance from a Pygments style dictionary.
+ (One that maps Token objects to style strings.)
+ """
+ pygments_style = []
+
+ for token, style in pygments_dict.items():
+ pygments_style.append((pygments_token_to_classname(token), style))
+
+ return Style(pygments_style)
+
+
+def pygments_token_to_classname(token: Token) -> str:
+ """
+ Turn e.g. `Token.Name.Exception` into `'pygments.name.exception'`.
+
+ (Our Pygments lexer will also turn the tokens that pygments produces in a
+ prompt_toolkit list of fragments that match these styling rules.)
+ """
+ parts = ("pygments",) + token
+ return ".".join(parts).lower()
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/style.py b/lib/python3.12/site-packages/prompt_toolkit/styles/style.py
new file mode 100644
index 0000000000000000000000000000000000000000..8aab8e1b21a2339eef88dfb01ab3b8f82083c068
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/style.py
@@ -0,0 +1,407 @@
+"""
+Tool for creating styles from a dictionary.
+"""
+
+from __future__ import annotations
+
+import itertools
+import re
+from enum import Enum
+from typing import Hashable, TypeVar
+
+from prompt_toolkit.cache import SimpleCache
+
+from .base import (
+ ANSI_COLOR_NAMES,
+ ANSI_COLOR_NAMES_ALIASES,
+ DEFAULT_ATTRS,
+ Attrs,
+ BaseStyle,
+)
+from .named_colors import NAMED_COLORS
+
+__all__ = [
+ "Style",
+ "parse_color",
+ "Priority",
+ "merge_styles",
+]
+
+_named_colors_lowercase = {k.lower(): v.lstrip("#") for k, v in NAMED_COLORS.items()}
+
+
+def parse_color(text: str) -> str:
+ """
+ Parse/validate color format.
+
+ Like in Pygments, but also support the ANSI color names.
+ (These will map to the colors of the 16 color palette.)
+ """
+ # ANSI color names.
+ if text in ANSI_COLOR_NAMES:
+ return text
+ if text in ANSI_COLOR_NAMES_ALIASES:
+ return ANSI_COLOR_NAMES_ALIASES[text]
+
+ # 140 named colors.
+ try:
+ # Replace by 'hex' value.
+ return _named_colors_lowercase[text.lower()]
+ except KeyError:
+ pass
+
+ # Hex codes.
+ if text[0:1] == "#":
+ col = text[1:]
+
+ # Keep this for backwards-compatibility (Pygments does it).
+ # I don't like the '#' prefix for named colors.
+ if col in ANSI_COLOR_NAMES:
+ return col
+ elif col in ANSI_COLOR_NAMES_ALIASES:
+ return ANSI_COLOR_NAMES_ALIASES[col]
+
+ # 6 digit hex color.
+ elif len(col) == 6:
+ return col
+
+ # 3 digit hex color.
+ elif len(col) == 3:
+ return col[0] * 2 + col[1] * 2 + col[2] * 2
+
+ # Default.
+ elif text in ("", "default"):
+ return text
+
+ raise ValueError(f"Wrong color format {text!r}")
+
+
+# Attributes, when they are not filled in by a style. None means that we take
+# the value from the parent.
+_EMPTY_ATTRS = Attrs(
+ color=None,
+ bgcolor=None,
+ bold=None,
+ underline=None,
+ strike=None,
+ italic=None,
+ blink=None,
+ reverse=None,
+ hidden=None,
+ dim=None,
+)
+
+
+def _expand_classname(classname: str) -> list[str]:
+ """
+ Split a single class name at the `.` operator, and build a list of classes.
+
+ E.g. 'a.b.c' becomes ['a', 'a.b', 'a.b.c']
+ """
+ result = []
+ parts = classname.split(".")
+
+ for i in range(1, len(parts) + 1):
+ result.append(".".join(parts[:i]).lower())
+
+ return result
+
+
+def _parse_style_str(style_str: str) -> Attrs:
+ """
+ Take a style string, e.g. 'bg:red #88ff00 class:title'
+ and return a `Attrs` instance.
+ """
+ # Start from default Attrs.
+ if "noinherit" in style_str:
+ attrs = DEFAULT_ATTRS
+ else:
+ attrs = _EMPTY_ATTRS
+
+ # Now update with the given attributes.
+ for part in style_str.split():
+ if part == "noinherit":
+ pass
+ elif part == "bold":
+ attrs = attrs._replace(bold=True)
+ elif part == "nobold":
+ attrs = attrs._replace(bold=False)
+ elif part == "italic":
+ attrs = attrs._replace(italic=True)
+ elif part == "noitalic":
+ attrs = attrs._replace(italic=False)
+ elif part == "underline":
+ attrs = attrs._replace(underline=True)
+ elif part == "nounderline":
+ attrs = attrs._replace(underline=False)
+ elif part == "strike":
+ attrs = attrs._replace(strike=True)
+ elif part == "nostrike":
+ attrs = attrs._replace(strike=False)
+
+ # prompt_toolkit extensions. Not in Pygments.
+ elif part == "blink":
+ attrs = attrs._replace(blink=True)
+ elif part == "noblink":
+ attrs = attrs._replace(blink=False)
+ elif part == "reverse":
+ attrs = attrs._replace(reverse=True)
+ elif part == "noreverse":
+ attrs = attrs._replace(reverse=False)
+ elif part == "hidden":
+ attrs = attrs._replace(hidden=True)
+ elif part == "nohidden":
+ attrs = attrs._replace(hidden=False)
+ elif part == "dim":
+ attrs = attrs._replace(dim=True)
+ elif part == "nodim":
+ attrs = attrs._replace(dim=False)
+
+ # Pygments properties that we ignore.
+ elif part in ("roman", "sans", "mono"):
+ pass
+ elif part.startswith("border:"):
+ pass
+
+ # Ignore pieces in between square brackets. This is internal stuff.
+ # Like '[transparent]' or '[set-cursor-position]'.
+ elif part.startswith("[") and part.endswith("]"):
+ pass
+
+ # Colors.
+ elif part.startswith("bg:"):
+ attrs = attrs._replace(bgcolor=parse_color(part[3:]))
+ elif part.startswith("fg:"): # The 'fg:' prefix is optional.
+ attrs = attrs._replace(color=parse_color(part[3:]))
+ else:
+ attrs = attrs._replace(color=parse_color(part))
+
+ return attrs
+
+
+CLASS_NAMES_RE = re.compile(r"^[a-z0-9.\s_-]*$") # This one can't contain a comma!
+
+
+class Priority(Enum):
+ """
+ The priority of the rules, when a style is created from a dictionary.
+
+ In a `Style`, rules that are defined later will always override previous
+ defined rules, however in a dictionary, the key order was arbitrary before
+ Python 3.6. This means that the style could change at random between rules.
+
+ We have two options:
+
+ - `DICT_KEY_ORDER`: This means, iterate through the dictionary, and take
+ the key/value pairs in order as they come. This is a good option if you
+ have Python >3.6. Rules at the end will override rules at the beginning.
+ - `MOST_PRECISE`: keys that are defined with most precision will get higher
+ priority. (More precise means: more elements.)
+ """
+
+ DICT_KEY_ORDER = "KEY_ORDER"
+ MOST_PRECISE = "MOST_PRECISE"
+
+
+# We don't support Python versions older than 3.6 anymore, so we can always
+# depend on dictionary ordering. This is the default.
+default_priority = Priority.DICT_KEY_ORDER
+
+
+class Style(BaseStyle):
+ """
+ Create a ``Style`` instance from a list of style rules.
+
+ The `style_rules` is supposed to be a list of ('classnames', 'style') tuples.
+ The classnames are a whitespace separated string of class names and the
+ style string is just like a Pygments style definition, but with a few
+ additions: it supports 'reverse' and 'blink'.
+
+ Later rules always override previous rules.
+
+ Usage::
+
+ Style([
+ ('title', '#ff0000 bold underline'),
+ ('something-else', 'reverse'),
+ ('class1 class2', 'reverse'),
+ ])
+
+ The ``from_dict`` classmethod is similar, but takes a dictionary as input.
+ """
+
+ def __init__(self, style_rules: list[tuple[str, str]]) -> None:
+ class_names_and_attrs = []
+
+ # Loop through the rules in the order they were defined.
+ # Rules that are defined later get priority.
+ for class_names, style_str in style_rules:
+ assert CLASS_NAMES_RE.match(class_names), repr(class_names)
+
+ # The order of the class names doesn't matter.
+ # (But the order of rules does matter.)
+ class_names_set = frozenset(class_names.lower().split())
+ attrs = _parse_style_str(style_str)
+
+ class_names_and_attrs.append((class_names_set, attrs))
+
+ self._style_rules = style_rules
+ self.class_names_and_attrs = class_names_and_attrs
+
+ @property
+ def style_rules(self) -> list[tuple[str, str]]:
+ return self._style_rules
+
+ @classmethod
+ def from_dict(
+ cls, style_dict: dict[str, str], priority: Priority = default_priority
+ ) -> Style:
+ """
+ :param style_dict: Style dictionary.
+ :param priority: `Priority` value.
+ """
+ if priority == Priority.MOST_PRECISE:
+
+ def key(item: tuple[str, str]) -> int:
+ # Split on '.' and whitespace. Count elements.
+ return sum(len(i.split(".")) for i in item[0].split())
+
+ return cls(sorted(style_dict.items(), key=key))
+ else:
+ return cls(list(style_dict.items()))
+
+ def get_attrs_for_style_str(
+ self, style_str: str, default: Attrs = DEFAULT_ATTRS
+ ) -> Attrs:
+ """
+ Get `Attrs` for the given style string.
+ """
+ list_of_attrs = [default]
+ class_names: set[str] = set()
+
+ # Apply default styling.
+ for names, attr in self.class_names_and_attrs:
+ if not names:
+ list_of_attrs.append(attr)
+
+ # Go from left to right through the style string. Things on the right
+ # take precedence.
+ for part in style_str.split():
+ # This part represents a class.
+ # Do lookup of this class name in the style definition, as well
+ # as all class combinations that we have so far.
+ if part.startswith("class:"):
+ # Expand all class names (comma separated list).
+ new_class_names = []
+ for p in part[6:].lower().split(","):
+ new_class_names.extend(_expand_classname(p))
+
+ for new_name in new_class_names:
+ # Build a set of all possible class combinations to be applied.
+ combos = set()
+ combos.add(frozenset([new_name]))
+
+ for count in range(1, len(class_names) + 1):
+ for c2 in itertools.combinations(class_names, count):
+ combos.add(frozenset(c2 + (new_name,)))
+
+ # Apply the styles that match these class names.
+ for names, attr in self.class_names_and_attrs:
+ if names in combos:
+ list_of_attrs.append(attr)
+
+ class_names.add(new_name)
+
+ # Process inline style.
+ else:
+ inline_attrs = _parse_style_str(part)
+ list_of_attrs.append(inline_attrs)
+
+ return _merge_attrs(list_of_attrs)
+
+ def invalidation_hash(self) -> Hashable:
+ return id(self.class_names_and_attrs)
+
+
+_T = TypeVar("_T")
+
+
+def _merge_attrs(list_of_attrs: list[Attrs]) -> Attrs:
+ """
+ Take a list of :class:`.Attrs` instances and merge them into one.
+ Every `Attr` in the list can override the styling of the previous one. So,
+ the last one has highest priority.
+ """
+
+ def _or(*values: _T) -> _T:
+ "Take first not-None value, starting at the end."
+ for v in values[::-1]:
+ if v is not None:
+ return v
+ raise ValueError # Should not happen, there's always one non-null value.
+
+ return Attrs(
+ color=_or("", *[a.color for a in list_of_attrs]),
+ bgcolor=_or("", *[a.bgcolor for a in list_of_attrs]),
+ bold=_or(False, *[a.bold for a in list_of_attrs]),
+ underline=_or(False, *[a.underline for a in list_of_attrs]),
+ strike=_or(False, *[a.strike for a in list_of_attrs]),
+ italic=_or(False, *[a.italic for a in list_of_attrs]),
+ blink=_or(False, *[a.blink for a in list_of_attrs]),
+ reverse=_or(False, *[a.reverse for a in list_of_attrs]),
+ hidden=_or(False, *[a.hidden for a in list_of_attrs]),
+ dim=_or(False, *[a.dim for a in list_of_attrs]),
+ )
+
+
+def merge_styles(styles: list[BaseStyle]) -> _MergedStyle:
+ """
+ Merge multiple `Style` objects.
+ """
+ styles = [s for s in styles if s is not None]
+ return _MergedStyle(styles)
+
+
+class _MergedStyle(BaseStyle):
+ """
+ Merge multiple `Style` objects into one.
+ This is supposed to ensure consistency: if any of the given styles changes,
+ then this style will be updated.
+ """
+
+ # NOTE: previously, we used an algorithm where we did not generate the
+ # combined style. Instead this was a proxy that called one style
+ # after the other, passing the outcome of the previous style as the
+ # default for the next one. This did not work, because that way, the
+ # priorities like described in the `Style` class don't work.
+ # 'class:aborted' was for instance never displayed in gray, because
+ # the next style specified a default color for any text. (The
+ # explicit styling of class:aborted should have taken priority,
+ # because it was more precise.)
+ def __init__(self, styles: list[BaseStyle]) -> None:
+ self.styles = styles
+ self._style: SimpleCache[Hashable, Style] = SimpleCache(maxsize=1)
+
+ @property
+ def _merged_style(self) -> Style:
+ "The `Style` object that has the other styles merged together."
+
+ def get() -> Style:
+ return Style(self.style_rules)
+
+ return self._style.get(self.invalidation_hash(), get)
+
+ @property
+ def style_rules(self) -> list[tuple[str, str]]:
+ style_rules = []
+ for s in self.styles:
+ style_rules.extend(s.style_rules)
+ return style_rules
+
+ def get_attrs_for_style_str(
+ self, style_str: str, default: Attrs = DEFAULT_ATTRS
+ ) -> Attrs:
+ return self._merged_style.get_attrs_for_style_str(style_str, default)
+
+ def invalidation_hash(self) -> Hashable:
+ return tuple(s.invalidation_hash() for s in self.styles)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/styles/style_transformation.py b/lib/python3.12/site-packages/prompt_toolkit/styles/style_transformation.py
new file mode 100644
index 0000000000000000000000000000000000000000..e8d5b0a89fe52a4f0e96e4369a761e4a2695127f
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/styles/style_transformation.py
@@ -0,0 +1,374 @@
+"""
+Collection of style transformations.
+
+Think of it as a kind of color post processing after the rendering is done.
+This could be used for instance to change the contrast/saturation; swap light
+and dark colors or even change certain colors for other colors.
+
+When the UI is rendered, these transformations can be applied right after the
+style strings are turned into `Attrs` objects that represent the actual
+formatting.
+"""
+
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from colorsys import hls_to_rgb, rgb_to_hls
+from typing import Callable, Hashable, Sequence
+
+from prompt_toolkit.cache import memoized
+from prompt_toolkit.filters import FilterOrBool, to_filter
+from prompt_toolkit.utils import AnyFloat, to_float, to_str
+
+from .base import ANSI_COLOR_NAMES, Attrs
+from .style import parse_color
+
+__all__ = [
+ "StyleTransformation",
+ "SwapLightAndDarkStyleTransformation",
+ "ReverseStyleTransformation",
+ "SetDefaultColorStyleTransformation",
+ "AdjustBrightnessStyleTransformation",
+ "DummyStyleTransformation",
+ "ConditionalStyleTransformation",
+ "DynamicStyleTransformation",
+ "merge_style_transformations",
+]
+
+
+class StyleTransformation(metaclass=ABCMeta):
+ """
+ Base class for any style transformation.
+ """
+
+ @abstractmethod
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ """
+ Take an `Attrs` object and return a new `Attrs` object.
+
+ Remember that the color formats can be either "ansi..." or a 6 digit
+ lowercase hexadecimal color (without '#' prefix).
+ """
+
+ def invalidation_hash(self) -> Hashable:
+ """
+ When this changes, the cache should be invalidated.
+ """
+ return f"{self.__class__.__name__}-{id(self)}"
+
+
+class SwapLightAndDarkStyleTransformation(StyleTransformation):
+ """
+ Turn dark colors into light colors and the other way around.
+
+ This is meant to make color schemes that work on a dark background usable
+ on a light background (and the other way around).
+
+ Notice that this doesn't swap foreground and background like "reverse"
+ does. It turns light green into dark green and the other way around.
+ Foreground and background colors are considered individually.
+
+ Also notice that when is used somewhere and no colors are given
+ in particular (like what is the default for the bottom toolbar), then this
+ doesn't change anything. This is what makes sense, because when the
+ 'default' color is chosen, it's what works best for the terminal, and
+ reverse works good with that.
+ """
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ """
+ Return the `Attrs` used when opposite luminosity should be used.
+ """
+ # Reverse colors.
+ attrs = attrs._replace(color=get_opposite_color(attrs.color))
+ attrs = attrs._replace(bgcolor=get_opposite_color(attrs.bgcolor))
+
+ return attrs
+
+
+class ReverseStyleTransformation(StyleTransformation):
+ """
+ Swap the 'reverse' attribute.
+
+ (This is still experimental.)
+ """
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ return attrs._replace(reverse=not attrs.reverse)
+
+
+class SetDefaultColorStyleTransformation(StyleTransformation):
+ """
+ Set default foreground/background color for output that doesn't specify
+ anything. This is useful for overriding the terminal default colors.
+
+ :param fg: Color string or callable that returns a color string for the
+ foreground.
+ :param bg: Like `fg`, but for the background.
+ """
+
+ def __init__(
+ self, fg: str | Callable[[], str], bg: str | Callable[[], str]
+ ) -> None:
+ self.fg = fg
+ self.bg = bg
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ if attrs.bgcolor in ("", "default"):
+ attrs = attrs._replace(bgcolor=parse_color(to_str(self.bg)))
+
+ if attrs.color in ("", "default"):
+ attrs = attrs._replace(color=parse_color(to_str(self.fg)))
+
+ return attrs
+
+ def invalidation_hash(self) -> Hashable:
+ return (
+ "set-default-color",
+ to_str(self.fg),
+ to_str(self.bg),
+ )
+
+
+class AdjustBrightnessStyleTransformation(StyleTransformation):
+ """
+ Adjust the brightness to improve the rendering on either dark or light
+ backgrounds.
+
+ For dark backgrounds, it's best to increase `min_brightness`. For light
+ backgrounds it's best to decrease `max_brightness`. Usually, only one
+ setting is adjusted.
+
+ This will only change the brightness for text that has a foreground color
+ defined, but no background color. It works best for 256 or true color
+ output.
+
+ .. note:: Notice that there is no universal way to detect whether the
+ application is running in a light or dark terminal. As a
+ developer of an command line application, you'll have to make
+ this configurable for the user.
+
+ :param min_brightness: Float between 0.0 and 1.0 or a callable that returns
+ a float.
+ :param max_brightness: Float between 0.0 and 1.0 or a callable that returns
+ a float.
+ """
+
+ def __init__(
+ self, min_brightness: AnyFloat = 0.0, max_brightness: AnyFloat = 1.0
+ ) -> None:
+ self.min_brightness = min_brightness
+ self.max_brightness = max_brightness
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ min_brightness = to_float(self.min_brightness)
+ max_brightness = to_float(self.max_brightness)
+ assert 0 <= min_brightness <= 1
+ assert 0 <= max_brightness <= 1
+
+ # Don't do anything if the whole brightness range is acceptable.
+ # This also avoids turning ansi colors into RGB sequences.
+ if min_brightness == 0.0 and max_brightness == 1.0:
+ return attrs
+
+ # If a foreground color is given without a background color.
+ no_background = not attrs.bgcolor or attrs.bgcolor == "default"
+ has_fgcolor = attrs.color and attrs.color != "ansidefault"
+
+ if has_fgcolor and no_background:
+ # Calculate new RGB values.
+ r, g, b = self._color_to_rgb(attrs.color or "")
+ hue, brightness, saturation = rgb_to_hls(r, g, b)
+ brightness = self._interpolate_brightness(
+ brightness, min_brightness, max_brightness
+ )
+ r, g, b = hls_to_rgb(hue, brightness, saturation)
+ new_color = f"{int(r * 255):02x}{int(g * 255):02x}{int(b * 255):02x}"
+
+ attrs = attrs._replace(color=new_color)
+
+ return attrs
+
+ def _color_to_rgb(self, color: str) -> tuple[float, float, float]:
+ """
+ Parse `style.Attrs` color into RGB tuple.
+ """
+ # Do RGB lookup for ANSI colors.
+ try:
+ from prompt_toolkit.output.vt100 import ANSI_COLORS_TO_RGB
+
+ r, g, b = ANSI_COLORS_TO_RGB[color]
+ return r / 255.0, g / 255.0, b / 255.0
+ except KeyError:
+ pass
+
+ # Parse RRGGBB format.
+ return (
+ int(color[0:2], 16) / 255.0,
+ int(color[2:4], 16) / 255.0,
+ int(color[4:6], 16) / 255.0,
+ )
+
+ # NOTE: we don't have to support named colors here. They are already
+ # transformed into RGB values in `style.parse_color`.
+
+ def _interpolate_brightness(
+ self, value: float, min_brightness: float, max_brightness: float
+ ) -> float:
+ """
+ Map the brightness to the (min_brightness..max_brightness) range.
+ """
+ return min_brightness + (max_brightness - min_brightness) * value
+
+ def invalidation_hash(self) -> Hashable:
+ return (
+ "adjust-brightness",
+ to_float(self.min_brightness),
+ to_float(self.max_brightness),
+ )
+
+
+class DummyStyleTransformation(StyleTransformation):
+ """
+ Don't transform anything at all.
+ """
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ return attrs
+
+ def invalidation_hash(self) -> Hashable:
+ # Always return the same hash for these dummy instances.
+ return "dummy-style-transformation"
+
+
+class DynamicStyleTransformation(StyleTransformation):
+ """
+ StyleTransformation class that can dynamically returns any
+ `StyleTransformation`.
+
+ :param get_style_transformation: Callable that returns a
+ :class:`.StyleTransformation` instance.
+ """
+
+ def __init__(
+ self, get_style_transformation: Callable[[], StyleTransformation | None]
+ ) -> None:
+ self.get_style_transformation = get_style_transformation
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ style_transformation = (
+ self.get_style_transformation() or DummyStyleTransformation()
+ )
+ return style_transformation.transform_attrs(attrs)
+
+ def invalidation_hash(self) -> Hashable:
+ style_transformation = (
+ self.get_style_transformation() or DummyStyleTransformation()
+ )
+ return style_transformation.invalidation_hash()
+
+
+class ConditionalStyleTransformation(StyleTransformation):
+ """
+ Apply the style transformation depending on a condition.
+ """
+
+ def __init__(
+ self, style_transformation: StyleTransformation, filter: FilterOrBool
+ ) -> None:
+ self.style_transformation = style_transformation
+ self.filter = to_filter(filter)
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ if self.filter():
+ return self.style_transformation.transform_attrs(attrs)
+ return attrs
+
+ def invalidation_hash(self) -> Hashable:
+ return (self.filter(), self.style_transformation.invalidation_hash())
+
+
+class _MergedStyleTransformation(StyleTransformation):
+ def __init__(self, style_transformations: Sequence[StyleTransformation]) -> None:
+ self.style_transformations = style_transformations
+
+ def transform_attrs(self, attrs: Attrs) -> Attrs:
+ for transformation in self.style_transformations:
+ attrs = transformation.transform_attrs(attrs)
+ return attrs
+
+ def invalidation_hash(self) -> Hashable:
+ return tuple(t.invalidation_hash() for t in self.style_transformations)
+
+
+def merge_style_transformations(
+ style_transformations: Sequence[StyleTransformation],
+) -> StyleTransformation:
+ """
+ Merge multiple transformations together.
+ """
+ return _MergedStyleTransformation(style_transformations)
+
+
+# Dictionary that maps ANSI color names to their opposite. This is useful for
+# turning color schemes that are optimized for a black background usable for a
+# white background.
+OPPOSITE_ANSI_COLOR_NAMES = {
+ "ansidefault": "ansidefault",
+ "ansiblack": "ansiwhite",
+ "ansired": "ansibrightred",
+ "ansigreen": "ansibrightgreen",
+ "ansiyellow": "ansibrightyellow",
+ "ansiblue": "ansibrightblue",
+ "ansimagenta": "ansibrightmagenta",
+ "ansicyan": "ansibrightcyan",
+ "ansigray": "ansibrightblack",
+ "ansiwhite": "ansiblack",
+ "ansibrightred": "ansired",
+ "ansibrightgreen": "ansigreen",
+ "ansibrightyellow": "ansiyellow",
+ "ansibrightblue": "ansiblue",
+ "ansibrightmagenta": "ansimagenta",
+ "ansibrightcyan": "ansicyan",
+ "ansibrightblack": "ansigray",
+}
+assert set(OPPOSITE_ANSI_COLOR_NAMES.keys()) == set(ANSI_COLOR_NAMES)
+assert set(OPPOSITE_ANSI_COLOR_NAMES.values()) == set(ANSI_COLOR_NAMES)
+
+
+@memoized()
+def get_opposite_color(colorname: str | None) -> str | None:
+ """
+ Take a color name in either 'ansi...' format or 6 digit RGB, return the
+ color of opposite luminosity (same hue/saturation).
+
+ This is used for turning color schemes that work on a light background
+ usable on a dark background.
+ """
+ if colorname is None: # Because color/bgcolor can be None in `Attrs`.
+ return None
+
+ # Special values.
+ if colorname in ("", "default"):
+ return colorname
+
+ # Try ANSI color names.
+ try:
+ return OPPOSITE_ANSI_COLOR_NAMES[colorname]
+ except KeyError:
+ # Try 6 digit RGB colors.
+ r = int(colorname[:2], 16) / 255.0
+ g = int(colorname[2:4], 16) / 255.0
+ b = int(colorname[4:6], 16) / 255.0
+
+ h, l, s = rgb_to_hls(r, g, b)
+
+ l = 1 - l
+
+ r, g, b = hls_to_rgb(h, l, s)
+
+ r = int(r * 255)
+ g = int(g * 255)
+ b = int(b * 255)
+
+ return f"{r:02x}{g:02x}{b:02x}"
diff --git a/lib/python3.12/site-packages/prompt_toolkit/token.py b/lib/python3.12/site-packages/prompt_toolkit/token.py
new file mode 100644
index 0000000000000000000000000000000000000000..e97893de673dcb1c761bcaea0b941d1666a87bc0
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/token.py
@@ -0,0 +1,9 @@
+""" """
+
+from __future__ import annotations
+
+__all__ = [
+ "ZeroWidthEscape",
+]
+
+ZeroWidthEscape = "[ZeroWidthEscape]"
diff --git a/lib/python3.12/site-packages/prompt_toolkit/utils.py b/lib/python3.12/site-packages/prompt_toolkit/utils.py
new file mode 100644
index 0000000000000000000000000000000000000000..1a99a2868a759f4bcff5abea34eac87abcef1e84
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/utils.py
@@ -0,0 +1,327 @@
+from __future__ import annotations
+
+import os
+import signal
+import sys
+import threading
+from collections import deque
+from typing import (
+ Callable,
+ ContextManager,
+ Dict,
+ Generator,
+ Generic,
+ TypeVar,
+ Union,
+)
+
+from wcwidth import wcwidth
+
+__all__ = [
+ "Event",
+ "DummyContext",
+ "get_cwidth",
+ "suspend_to_background_supported",
+ "is_conemu_ansi",
+ "is_windows",
+ "in_main_thread",
+ "get_bell_environment_variable",
+ "get_term_environment_variable",
+ "take_using_weights",
+ "to_str",
+ "to_int",
+ "AnyFloat",
+ "to_float",
+ "is_dumb_terminal",
+]
+
+# Used to ensure sphinx autodoc does not try to import platform-specific
+# stuff when documenting win32.py modules.
+SPHINX_AUTODOC_RUNNING = "sphinx.ext.autodoc" in sys.modules
+
+_Sender = TypeVar("_Sender", covariant=True)
+
+
+class Event(Generic[_Sender]):
+ """
+ Simple event to which event handlers can be attached. For instance::
+
+ class Cls:
+ def __init__(self):
+ # Define event. The first parameter is the sender.
+ self.event = Event(self)
+
+ obj = Cls()
+
+ def handler(sender):
+ pass
+
+ # Add event handler by using the += operator.
+ obj.event += handler
+
+ # Fire event.
+ obj.event()
+ """
+
+ def __init__(
+ self, sender: _Sender, handler: Callable[[_Sender], None] | None = None
+ ) -> None:
+ self.sender = sender
+ self._handlers: list[Callable[[_Sender], None]] = []
+
+ if handler is not None:
+ self += handler
+
+ def __call__(self) -> None:
+ "Fire event."
+ for handler in self._handlers:
+ handler(self.sender)
+
+ def fire(self) -> None:
+ "Alias for just calling the event."
+ self()
+
+ def add_handler(self, handler: Callable[[_Sender], None]) -> None:
+ """
+ Add another handler to this callback.
+ (Handler should be a callable that takes exactly one parameter: the
+ sender object.)
+ """
+ # Add to list of event handlers.
+ self._handlers.append(handler)
+
+ def remove_handler(self, handler: Callable[[_Sender], None]) -> None:
+ """
+ Remove a handler from this callback.
+ """
+ if handler in self._handlers:
+ self._handlers.remove(handler)
+
+ def __iadd__(self, handler: Callable[[_Sender], None]) -> Event[_Sender]:
+ """
+ `event += handler` notation for adding a handler.
+ """
+ self.add_handler(handler)
+ return self
+
+ def __isub__(self, handler: Callable[[_Sender], None]) -> Event[_Sender]:
+ """
+ `event -= handler` notation for removing a handler.
+ """
+ self.remove_handler(handler)
+ return self
+
+
+class DummyContext(ContextManager[None]):
+ """
+ (contextlib.nested is not available on Py3)
+ """
+
+ def __enter__(self) -> None:
+ pass
+
+ def __exit__(self, *a: object) -> None:
+ pass
+
+
+class _CharSizesCache(Dict[str, int]):
+ """
+ Cache for wcwidth sizes.
+ """
+
+ LONG_STRING_MIN_LEN = 64 # Minimum string length for considering it long.
+ MAX_LONG_STRINGS = 16 # Maximum number of long strings to remember.
+
+ def __init__(self) -> None:
+ super().__init__()
+ # Keep track of the "long" strings in this cache.
+ self._long_strings: deque[str] = deque()
+
+ def __missing__(self, string: str) -> int:
+ # Note: We use the `max(0, ...` because some non printable control
+ # characters, like e.g. Ctrl-underscore get a -1 wcwidth value.
+ # It can be possible that these characters end up in the input
+ # text.
+ result: int
+ if len(string) == 1:
+ result = max(0, wcwidth(string))
+ else:
+ result = sum(self[c] for c in string)
+
+ # Store in cache.
+ self[string] = result
+
+ # Rotate long strings.
+ # (It's hard to tell what we can consider short...)
+ if len(string) > self.LONG_STRING_MIN_LEN:
+ long_strings = self._long_strings
+ long_strings.append(string)
+
+ if len(long_strings) > self.MAX_LONG_STRINGS:
+ key_to_remove = long_strings.popleft()
+ if key_to_remove in self:
+ del self[key_to_remove]
+
+ return result
+
+
+_CHAR_SIZES_CACHE = _CharSizesCache()
+
+
+def get_cwidth(string: str) -> int:
+ """
+ Return width of a string. Wrapper around ``wcwidth``.
+ """
+ return _CHAR_SIZES_CACHE[string]
+
+
+def suspend_to_background_supported() -> bool:
+ """
+ Returns `True` when the Python implementation supports
+ suspend-to-background. This is typically `False' on Windows systems.
+ """
+ return hasattr(signal, "SIGTSTP")
+
+
+def is_windows() -> bool:
+ """
+ True when we are using Windows.
+ """
+ return sys.platform == "win32" # Not 'darwin' or 'linux2'
+
+
+def is_windows_vt100_supported() -> bool:
+ """
+ True when we are using Windows, but VT100 escape sequences are supported.
+ """
+ if sys.platform == "win32":
+ # Import needs to be inline. Windows libraries are not always available.
+ from prompt_toolkit.output.windows10 import is_win_vt100_enabled
+
+ return is_win_vt100_enabled()
+
+ return False
+
+
+def is_conemu_ansi() -> bool:
+ """
+ True when the ConEmu Windows console is used.
+ """
+ return sys.platform == "win32" and os.environ.get("ConEmuANSI", "OFF") == "ON"
+
+
+def in_main_thread() -> bool:
+ """
+ True when the current thread is the main thread.
+ """
+ return threading.current_thread().__class__.__name__ == "_MainThread"
+
+
+def get_bell_environment_variable() -> bool:
+ """
+ True if env variable is set to true (true, TRUE, True, 1).
+ """
+ value = os.environ.get("PROMPT_TOOLKIT_BELL", "true")
+ return value.lower() in ("1", "true")
+
+
+def get_term_environment_variable() -> str:
+ "Return the $TERM environment variable."
+ return os.environ.get("TERM", "")
+
+
+_T = TypeVar("_T")
+
+
+def take_using_weights(
+ items: list[_T], weights: list[int]
+) -> Generator[_T, None, None]:
+ """
+ Generator that keeps yielding items from the items list, in proportion to
+ their weight. For instance::
+
+ # Getting the first 70 items from this generator should have yielded 10
+ # times A, 20 times B and 40 times C, all distributed equally..
+ take_using_weights(['A', 'B', 'C'], [5, 10, 20])
+
+ :param items: List of items to take from.
+ :param weights: Integers representing the weight. (Numbers have to be
+ integers, not floats.)
+ """
+ assert len(items) == len(weights)
+ assert len(items) > 0
+
+ # Remove items with zero-weight.
+ items2 = []
+ weights2 = []
+ for item, w in zip(items, weights):
+ if w > 0:
+ items2.append(item)
+ weights2.append(w)
+
+ items = items2
+ weights = weights2
+
+ # Make sure that we have some items left.
+ if not items:
+ raise ValueError("Did't got any items with a positive weight.")
+
+ #
+ already_taken = [0 for i in items]
+ item_count = len(items)
+ max_weight = max(weights)
+
+ i = 0
+ while True:
+ # Each iteration of this loop, we fill up until by (total_weight/max_weight).
+ adding = True
+ while adding:
+ adding = False
+
+ for item_i, item, weight in zip(range(item_count), items, weights):
+ if already_taken[item_i] < i * weight / float(max_weight):
+ yield item
+ already_taken[item_i] += 1
+ adding = True
+
+ i += 1
+
+
+def to_str(value: Callable[[], str] | str) -> str:
+ "Turn callable or string into string."
+ if callable(value):
+ return to_str(value())
+ else:
+ return str(value)
+
+
+def to_int(value: Callable[[], int] | int) -> int:
+ "Turn callable or int into int."
+ if callable(value):
+ return to_int(value())
+ else:
+ return int(value)
+
+
+AnyFloat = Union[Callable[[], float], float]
+
+
+def to_float(value: AnyFloat) -> float:
+ "Turn callable or float into float."
+ if callable(value):
+ return to_float(value())
+ else:
+ return float(value)
+
+
+def is_dumb_terminal(term: str | None = None) -> bool:
+ """
+ True if this terminal type is considered "dumb".
+
+ If so, we should fall back to the simplest possible form of line editing,
+ without cursor positioning and color support.
+ """
+ if term is None:
+ return is_dumb_terminal(os.environ.get("TERM", ""))
+
+ return term.lower() in ["dumb", "unknown"]
diff --git a/lib/python3.12/site-packages/prompt_toolkit/validation.py b/lib/python3.12/site-packages/prompt_toolkit/validation.py
new file mode 100644
index 0000000000000000000000000000000000000000..2b35d1fcc91873a9b8242495301854b0cb81e6a4
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/validation.py
@@ -0,0 +1,192 @@
+"""
+Input validation for a `Buffer`.
+(Validators will be called before accepting input.)
+"""
+
+from __future__ import annotations
+
+from abc import ABCMeta, abstractmethod
+from typing import Callable
+
+from prompt_toolkit.eventloop import run_in_executor_with_context
+
+from .document import Document
+from .filters import FilterOrBool, to_filter
+
+__all__ = [
+ "ConditionalValidator",
+ "ValidationError",
+ "Validator",
+ "ThreadedValidator",
+ "DummyValidator",
+ "DynamicValidator",
+]
+
+
+class ValidationError(Exception):
+ """
+ Error raised by :meth:`.Validator.validate`.
+
+ :param cursor_position: The cursor position where the error occurred.
+ :param message: Text.
+ """
+
+ def __init__(self, cursor_position: int = 0, message: str = "") -> None:
+ super().__init__(message)
+ self.cursor_position = cursor_position
+ self.message = message
+
+ def __repr__(self) -> str:
+ return f"{self.__class__.__name__}(cursor_position={self.cursor_position!r}, message={self.message!r})"
+
+
+class Validator(metaclass=ABCMeta):
+ """
+ Abstract base class for an input validator.
+
+ A validator is typically created in one of the following two ways:
+
+ - Either by overriding this class and implementing the `validate` method.
+ - Or by passing a callable to `Validator.from_callable`.
+
+ If the validation takes some time and needs to happen in a background
+ thread, this can be wrapped in a :class:`.ThreadedValidator`.
+ """
+
+ @abstractmethod
+ def validate(self, document: Document) -> None:
+ """
+ Validate the input.
+ If invalid, this should raise a :class:`.ValidationError`.
+
+ :param document: :class:`~prompt_toolkit.document.Document` instance.
+ """
+ pass
+
+ async def validate_async(self, document: Document) -> None:
+ """
+ Return a `Future` which is set when the validation is ready.
+ This function can be overloaded in order to provide an asynchronous
+ implementation.
+ """
+ try:
+ self.validate(document)
+ except ValidationError:
+ raise
+
+ @classmethod
+ def from_callable(
+ cls,
+ validate_func: Callable[[str], bool],
+ error_message: str = "Invalid input",
+ move_cursor_to_end: bool = False,
+ ) -> Validator:
+ """
+ Create a validator from a simple validate callable. E.g.:
+
+ .. code:: python
+
+ def is_valid(text):
+ return text in ['hello', 'world']
+ Validator.from_callable(is_valid, error_message='Invalid input')
+
+ :param validate_func: Callable that takes the input string, and returns
+ `True` if the input is valid input.
+ :param error_message: Message to be displayed if the input is invalid.
+ :param move_cursor_to_end: Move the cursor to the end of the input, if
+ the input is invalid.
+ """
+ return _ValidatorFromCallable(validate_func, error_message, move_cursor_to_end)
+
+
+class _ValidatorFromCallable(Validator):
+ """
+ Validate input from a simple callable.
+ """
+
+ def __init__(
+ self, func: Callable[[str], bool], error_message: str, move_cursor_to_end: bool
+ ) -> None:
+ self.func = func
+ self.error_message = error_message
+ self.move_cursor_to_end = move_cursor_to_end
+
+ def __repr__(self) -> str:
+ return f"Validator.from_callable({self.func!r})"
+
+ def validate(self, document: Document) -> None:
+ if not self.func(document.text):
+ if self.move_cursor_to_end:
+ index = len(document.text)
+ else:
+ index = 0
+
+ raise ValidationError(cursor_position=index, message=self.error_message)
+
+
+class ThreadedValidator(Validator):
+ """
+ Wrapper that runs input validation in a thread.
+ (Use this to prevent the user interface from becoming unresponsive if the
+ input validation takes too much time.)
+ """
+
+ def __init__(self, validator: Validator) -> None:
+ self.validator = validator
+
+ def validate(self, document: Document) -> None:
+ self.validator.validate(document)
+
+ async def validate_async(self, document: Document) -> None:
+ """
+ Run the `validate` function in a thread.
+ """
+
+ def run_validation_thread() -> None:
+ return self.validate(document)
+
+ await run_in_executor_with_context(run_validation_thread)
+
+
+class DummyValidator(Validator):
+ """
+ Validator class that accepts any input.
+ """
+
+ def validate(self, document: Document) -> None:
+ pass # Don't raise any exception.
+
+
+class ConditionalValidator(Validator):
+ """
+ Validator that can be switched on/off according to
+ a filter. (This wraps around another validator.)
+ """
+
+ def __init__(self, validator: Validator, filter: FilterOrBool) -> None:
+ self.validator = validator
+ self.filter = to_filter(filter)
+
+ def validate(self, document: Document) -> None:
+ # Call the validator only if the filter is active.
+ if self.filter():
+ self.validator.validate(document)
+
+
+class DynamicValidator(Validator):
+ """
+ Validator class that can dynamically returns any Validator.
+
+ :param get_validator: Callable that returns a :class:`.Validator` instance.
+ """
+
+ def __init__(self, get_validator: Callable[[], Validator | None]) -> None:
+ self.get_validator = get_validator
+
+ def validate(self, document: Document) -> None:
+ validator = self.get_validator() or DummyValidator()
+ validator.validate(document)
+
+ async def validate_async(self, document: Document) -> None:
+ validator = self.get_validator() or DummyValidator()
+ await validator.validate_async(document)
diff --git a/lib/python3.12/site-packages/prompt_toolkit/win32_types.py b/lib/python3.12/site-packages/prompt_toolkit/win32_types.py
new file mode 100644
index 0000000000000000000000000000000000000000..79283b8b06263808994236a11d7353e664eba486
--- /dev/null
+++ b/lib/python3.12/site-packages/prompt_toolkit/win32_types.py
@@ -0,0 +1,229 @@
+from __future__ import annotations
+
+from ctypes import Structure, Union, c_char, c_long, c_short, c_ulong
+from ctypes.wintypes import BOOL, DWORD, LPVOID, WCHAR, WORD
+from typing import TYPE_CHECKING
+
+# Input/Output standard device numbers. Note that these are not handle objects.
+# It's the `windll.kernel32.GetStdHandle` system call that turns them into a
+# real handle object.
+STD_INPUT_HANDLE = c_ulong(-10)
+STD_OUTPUT_HANDLE = c_ulong(-11)
+STD_ERROR_HANDLE = c_ulong(-12)
+
+
+class COORD(Structure):
+ """
+ Struct in wincon.h
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms682119(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ X: int
+ Y: int
+
+ _fields_ = [
+ ("X", c_short), # Short
+ ("Y", c_short), # Short
+ ]
+
+ def __repr__(self) -> str:
+ return "{}(X={!r}, Y={!r}, type_x={!r}, type_y={!r})".format(
+ self.__class__.__name__,
+ self.X,
+ self.Y,
+ type(self.X),
+ type(self.Y),
+ )
+
+
+class UNICODE_OR_ASCII(Union):
+ if TYPE_CHECKING:
+ AsciiChar: bytes
+ UnicodeChar: str
+
+ _fields_ = [
+ ("AsciiChar", c_char),
+ ("UnicodeChar", WCHAR),
+ ]
+
+
+class KEY_EVENT_RECORD(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms684166(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ KeyDown: int
+ RepeatCount: int
+ VirtualKeyCode: int
+ VirtualScanCode: int
+ uChar: UNICODE_OR_ASCII
+ ControlKeyState: int
+
+ _fields_ = [
+ ("KeyDown", c_long), # bool
+ ("RepeatCount", c_short), # word
+ ("VirtualKeyCode", c_short), # word
+ ("VirtualScanCode", c_short), # word
+ ("uChar", UNICODE_OR_ASCII), # Unicode or ASCII.
+ ("ControlKeyState", c_long), # double word
+ ]
+
+
+class MOUSE_EVENT_RECORD(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms684239(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ MousePosition: COORD
+ ButtonState: int
+ ControlKeyState: int
+ EventFlags: int
+
+ _fields_ = [
+ ("MousePosition", COORD),
+ ("ButtonState", c_long), # dword
+ ("ControlKeyState", c_long), # dword
+ ("EventFlags", c_long), # dword
+ ]
+
+
+class WINDOW_BUFFER_SIZE_RECORD(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms687093(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ Size: COORD
+
+ _fields_ = [("Size", COORD)]
+
+
+class MENU_EVENT_RECORD(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms684213(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ CommandId: int
+
+ _fields_ = [("CommandId", c_long)] # uint
+
+
+class FOCUS_EVENT_RECORD(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms683149(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ SetFocus: int
+
+ _fields_ = [("SetFocus", c_long)] # bool
+
+
+class EVENT_RECORD(Union):
+ if TYPE_CHECKING:
+ KeyEvent: KEY_EVENT_RECORD
+ MouseEvent: MOUSE_EVENT_RECORD
+ WindowBufferSizeEvent: WINDOW_BUFFER_SIZE_RECORD
+ MenuEvent: MENU_EVENT_RECORD
+ FocusEvent: FOCUS_EVENT_RECORD
+
+ _fields_ = [
+ ("KeyEvent", KEY_EVENT_RECORD),
+ ("MouseEvent", MOUSE_EVENT_RECORD),
+ ("WindowBufferSizeEvent", WINDOW_BUFFER_SIZE_RECORD),
+ ("MenuEvent", MENU_EVENT_RECORD),
+ ("FocusEvent", FOCUS_EVENT_RECORD),
+ ]
+
+
+class INPUT_RECORD(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/ms683499(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ EventType: int
+ Event: EVENT_RECORD
+
+ _fields_ = [("EventType", c_short), ("Event", EVENT_RECORD)] # word # Union.
+
+
+EventTypes = {
+ 1: "KeyEvent",
+ 2: "MouseEvent",
+ 4: "WindowBufferSizeEvent",
+ 8: "MenuEvent",
+ 16: "FocusEvent",
+}
+
+
+class SMALL_RECT(Structure):
+ """struct in wincon.h."""
+
+ if TYPE_CHECKING:
+ Left: int
+ Top: int
+ Right: int
+ Bottom: int
+
+ _fields_ = [
+ ("Left", c_short),
+ ("Top", c_short),
+ ("Right", c_short),
+ ("Bottom", c_short),
+ ]
+
+
+class CONSOLE_SCREEN_BUFFER_INFO(Structure):
+ """struct in wincon.h."""
+
+ if TYPE_CHECKING:
+ dwSize: COORD
+ dwCursorPosition: COORD
+ wAttributes: int
+ srWindow: SMALL_RECT
+ dwMaximumWindowSize: COORD
+
+ _fields_ = [
+ ("dwSize", COORD),
+ ("dwCursorPosition", COORD),
+ ("wAttributes", WORD),
+ ("srWindow", SMALL_RECT),
+ ("dwMaximumWindowSize", COORD),
+ ]
+
+ def __repr__(self) -> str:
+ return "CONSOLE_SCREEN_BUFFER_INFO({!r},{!r},{!r},{!r},{!r},{!r},{!r},{!r},{!r},{!r},{!r})".format(
+ self.dwSize.Y,
+ self.dwSize.X,
+ self.dwCursorPosition.Y,
+ self.dwCursorPosition.X,
+ self.wAttributes,
+ self.srWindow.Top,
+ self.srWindow.Left,
+ self.srWindow.Bottom,
+ self.srWindow.Right,
+ self.dwMaximumWindowSize.Y,
+ self.dwMaximumWindowSize.X,
+ )
+
+
+class SECURITY_ATTRIBUTES(Structure):
+ """
+ http://msdn.microsoft.com/en-us/library/windows/desktop/aa379560(v=vs.85).aspx
+ """
+
+ if TYPE_CHECKING:
+ nLength: int
+ lpSecurityDescriptor: int
+ bInheritHandle: int # BOOL comes back as 'int'.
+
+ _fields_ = [
+ ("nLength", DWORD),
+ ("lpSecurityDescriptor", LPVOID),
+ ("bInheritHandle", BOOL),
+ ]
diff --git a/lib/python3.12/site-packages/triton/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/triton/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..4caae506acbc2a01755f33fd758c41b319da6730
Binary files /dev/null and b/lib/python3.12/site-packages/triton/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/__pycache__/_internal_testing.cpython-312.pyc b/lib/python3.12/site-packages/triton/__pycache__/_internal_testing.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..49ed40d07711c936fdf07a7943168f24abc481ab
Binary files /dev/null and b/lib/python3.12/site-packages/triton/__pycache__/_internal_testing.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/__pycache__/errors.cpython-312.pyc b/lib/python3.12/site-packages/triton/__pycache__/errors.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..88422277fcd695f5d1ffc875c04537e6501918d3
Binary files /dev/null and b/lib/python3.12/site-packages/triton/__pycache__/errors.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/__pycache__/testing.cpython-312.pyc b/lib/python3.12/site-packages/triton/__pycache__/testing.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..4601916d48db12300e8d0f9a4f3476c94d4ed65a
Binary files /dev/null and b/lib/python3.12/site-packages/triton/__pycache__/testing.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/backends/__pycache__/__init__.cpython-312.pyc b/lib/python3.12/site-packages/triton/backends/__pycache__/__init__.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..a3c53ab657690ca6e86ef0d29ede537b68d13910
Binary files /dev/null and b/lib/python3.12/site-packages/triton/backends/__pycache__/__init__.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/backends/__pycache__/compiler.cpython-312.pyc b/lib/python3.12/site-packages/triton/backends/__pycache__/compiler.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..9e964d762466af44f1880b3299dbeffdde8b32b7
Binary files /dev/null and b/lib/python3.12/site-packages/triton/backends/__pycache__/compiler.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/backends/__pycache__/driver.cpython-312.pyc b/lib/python3.12/site-packages/triton/backends/__pycache__/driver.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..cd31fd38212b0cf7da30e7aa0aa70f9903080cfe
Binary files /dev/null and b/lib/python3.12/site-packages/triton/backends/__pycache__/driver.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/backends/amd/__pycache__/compiler.cpython-312.pyc b/lib/python3.12/site-packages/triton/backends/amd/__pycache__/compiler.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..be3931e542ad092cfd05976241f9efbfc5d812fb
Binary files /dev/null and b/lib/python3.12/site-packages/triton/backends/amd/__pycache__/compiler.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/backends/amd/__pycache__/driver.cpython-312.pyc b/lib/python3.12/site-packages/triton/backends/amd/__pycache__/driver.cpython-312.pyc
new file mode 100644
index 0000000000000000000000000000000000000000..87ee8d569f554d3b9be397b5c6c5e4662db75c51
Binary files /dev/null and b/lib/python3.12/site-packages/triton/backends/amd/__pycache__/driver.cpython-312.pyc differ
diff --git a/lib/python3.12/site-packages/triton/backends/amd/compiler.py b/lib/python3.12/site-packages/triton/backends/amd/compiler.py
new file mode 100644
index 0000000000000000000000000000000000000000..7a2ab1c53b2c4bf15c22f4403886b710c44ae6a5
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/compiler.py
@@ -0,0 +1,368 @@
+from triton.backends.compiler import BaseBackend, GPUTarget, AttrsDescriptor, register_descriptor
+from triton._C.libtriton import ir, passes, llvm, amd
+from dataclasses import dataclass
+from typing import Any, Dict, Tuple
+from types import ModuleType
+import hashlib
+import tempfile
+import os
+import re
+import subprocess
+import functools
+from pathlib import Path
+
+
+def min_dot_size(target: GPUTarget):
+ arch_str = target.arch
+ # CDNA 3.0 supports k==8 in all mfma variants except for int8
+ # (where the smallest `k` supported is 16)
+ if "gfx94" in arch_str:
+ return lambda lhsType, rhsType: (16, 16, 16) if (lhsType.is_int8() or rhsType.is_int8()) else (16, 16, 8)
+ # CDNA 2.0 always supports `k==8`
+ if "gfx9" in arch_str:
+ return lambda lhsType, rhsType: (16, 16, 8)
+ # Other architectures will only support 16,16,16
+ return lambda lhsType, rhsType: (16, 16, 16)
+
+
+@dataclass(frozen=True)
+class HIPOptions:
+ num_warps: int = 4
+ waves_per_eu: int = 1
+ num_stages: int = 2
+ num_ctas: int = 1
+ num_buffers_warp_spec: int = 0
+ num_consumer_groups: int = 0
+ reg_dec_producer: int = 0
+ reg_inc_consumer: int = 0
+ extern_libs: dict = None
+ cluster_dims: tuple = (1, 1, 1)
+ debug: bool = False
+ sanitize_overflow: bool = True
+ arch: str = None
+ supported_fp8_dtypes: Tuple[str] = ("fp8e5", )
+ deprecated_fp8_dtypes: Tuple[str] = ()
+ default_dot_input_precision: str = "ieee"
+ allowed_dot_input_precisions: Tuple[str] = ("ieee", )
+ enable_fp_fusion: bool = True
+ matrix_instr_nonkdim: int = 0
+ kpack: int = 1
+ allow_flush_denorm: bool = False
+ max_num_imprecise_acc_default: int = 0
+ backend_name: str = 'hip'
+
+ # The following option provides hints to the AMDGPU backend regarding instruction scheduling
+ # for all `tt.dot` operations in a kernel. The "default" variant preserves the default
+ # instruction scheduling of the AMDGPU backend which aims at maximizing occupancy.
+ # The option is experimental and may change at any time regarding its semantics and/or may
+ # be gone entirely anytime.
+ instruction_sched_variant: str = 'default'
+
+ def __post_init__(self):
+ default_libdir = Path(__file__).parent / 'lib'
+ extern_libs = {} if self.extern_libs is None else dict(self.extern_libs)
+ # Ignore user-defined warp size for gfx9
+ warp_size = 32 if 'gfx10' in self.arch or 'gfx11' in self.arch or 'gfx12' in self.arch else 64
+ object.__setattr__(self, 'warp_size', warp_size)
+ libs = ["ocml", "ockl"]
+ for lib in libs:
+ extern_libs[lib] = str(default_libdir / f'{lib}.bc')
+ object.__setattr__(self, 'extern_libs', tuple(extern_libs.items()))
+ assert self.num_warps > 0 and (self.num_warps & (self.num_warps - 1)) == 0, \
+ "num_warps must be a power of 2"
+
+ def hash(self):
+ key = '_'.join([f'{name}-{val}' for name, val in self.__dict__.items()])
+ return hashlib.sha256(key.encode("utf-8")).hexdigest()
+
+
+@register_descriptor
+class HIPAttrsDescriptor(AttrsDescriptor):
+ # This property asserts if the underlying storage area of a given pointer
+ # can be resepresented as a 32 bit integer. When this is true, we can be
+ # sure that all indices into the tensor behind that pointer can use 32-bit
+ # indexing. That opens the door for the AMD backend to use buffer load/store
+ # instrinsics, which requires this property. Buffer load/store intrinsics
+ # gives direct out-of-bound support and simplifies index calculation for
+ # lower register pressure.
+ __slots__ = ("pointer_range_32")
+
+ def _add_backend_properties(self, params=None, values=None):
+ self.property_values["tt.pointer_range"] = 32
+ if params is None or values is None:
+ return
+
+ self.arg_properties["tt.pointer_range"] = [
+ param.num for param, arg in zip(params, values) if HIPAttrsDescriptor.is_within2gb(arg)
+ and not param.do_not_specialize and not param.do_not_specialize_on_alignment
+ ]
+
+ @staticmethod
+ def is_within2gb(arg):
+ if hasattr(arg, "ptr_range"):
+ return arg.ptr_range() <= 2**31 - 1
+ if "torch.Tensor" in str(type(arg)) and hasattr(arg, "untyped_storage"):
+ # Please note that 2**31-1 is the max int32 positive limit
+ return arg.untyped_storage().size() <= 2**31 - 1
+ return False
+
+ @staticmethod
+ def get_property_key(val, align):
+ generic_key = AttrsDescriptor.get_property_key(val, align)
+ hip_key = "S" if HIPAttrsDescriptor.is_within2gb(val) else "N"
+ key = (generic_key + hip_key).replace("N", "")
+ return key if key else "N"
+
+
+class HIPBackend(BaseBackend):
+
+ @staticmethod
+ def supports_target(target: GPUTarget):
+ return target.backend == 'hip'
+
+ def __init__(self, target: GPUTarget) -> None:
+ super().__init__(target)
+ assert isinstance(target.arch, str)
+ self.binary_ext = "hsaco"
+
+ def parse_options(self, opts) -> Any:
+ args = {'arch': self.target.arch}
+
+ if "supported_fp8_dtypes" not in opts:
+ supported_fp8_dtypes = set(HIPOptions.supported_fp8_dtypes)
+ if self.target.arch in ('gfx940', 'gfx941', 'gfx942'):
+ supported_fp8_dtypes.update({'fp8e4b8', 'fp8e5b16'})
+ args["supported_fp8_dtypes"] = tuple(sorted(supported_fp8_dtypes))
+
+ if "enable_fp_fusion" not in opts:
+ args["enable_fp_fusion"] = os.getenv("TRITON_DEFAULT_FP_FUSION", "1") == "1"
+ args.update({k: opts[k] for k in HIPOptions.__dataclass_fields__.keys() if k in opts})
+ return HIPOptions(**args)
+
+ def pack_metadata(self, metadata):
+ return (
+ metadata.num_warps,
+ metadata.num_ctas,
+ metadata.shared,
+ metadata.cluster_dims[0],
+ metadata.cluster_dims[1],
+ metadata.cluster_dims[2],
+ )
+
+ def get_codegen_implementation(self):
+ codegen_fns = {"min_dot_size": min_dot_size(self.target)}
+ return codegen_fns
+
+ def get_module_map(self) -> Dict[str, ModuleType]:
+ from triton.language.extra.hip import libdevice
+ return {"triton.language.extra.libdevice": libdevice}
+
+ def load_dialects(self, ctx):
+ amd.load_dialects(ctx)
+
+ def get_attrs_descriptor(self, params, args):
+ return HIPAttrsDescriptor(params, args)
+
+ @staticmethod
+ def compute_spec_key(arg, align):
+ return HIPAttrsDescriptor.get_property_key(arg, align)
+
+ @staticmethod
+ def path_to_rocm_lld():
+ # Check env path for ld.lld
+ lld_env_path = os.getenv("TRITON_HIP_LLD_PATH")
+ if lld_env_path is not None:
+ lld = Path(lld_env_path)
+ if lld.is_file():
+ return lld
+ # Check backend for ld.lld (used for pytorch wheels)
+ lld = Path(__file__).parent / "llvm/bin/ld.lld"
+ if lld.is_file():
+ return lld
+ lld = Path("/opt/rocm/llvm/bin/ld.lld")
+ if lld.is_file():
+ return lld
+ lld = Path("/usr/bin/ld.lld")
+ if lld.is_file():
+ return lld
+ raise Exception("ROCm linker /opt/rocm/llvm/bin/ld.lld not found. Set 'TRITON_HIP_LLD_PATH' to its path.")
+
+ @staticmethod
+ def make_ttir(mod, metadata, options):
+ pm = ir.pass_manager(mod.context)
+ pm.enable_debug()
+ passes.common.add_inliner(pm)
+ passes.ttir.add_rewrite_tensor_pointer(pm)
+ passes.ttir.add_combine(pm)
+ passes.common.add_canonicalizer(pm)
+ passes.ttir.add_reorder_broadcast(pm)
+ passes.common.add_cse(pm)
+ passes.common.add_licm(pm)
+ passes.common.add_symbol_dce(pm)
+ passes.ttir.add_loop_unroll(pm)
+ pm.run(mod)
+ return mod
+
+ @staticmethod
+ def make_ttgir(mod, metadata, options):
+ pm = ir.pass_manager(mod.context)
+ pm.enable_debug()
+ passes.ttir.add_convert_to_ttgpuir(pm, f"hip:{options.arch}", options.num_warps, options.warp_size,
+ options.num_ctas)
+ pm.run(mod)
+ pm = ir.pass_manager(mod.context)
+ pm.enable_debug()
+ passes.ttgpuir.add_coalesce(pm)
+ passes.ttgpuir.add_remove_layout_conversions(pm)
+ passes.ttgpuir.add_optimize_thread_locality(pm)
+ amd.passes.ttgpuir.add_accelerate_matmul(pm, options.arch, options.matrix_instr_nonkdim, options.kpack)
+ passes.ttgpuir.add_remove_layout_conversions(pm)
+ amd.passes.ttgpuir.add_optimize_epilogue(pm)
+ passes.ttgpuir.add_optimize_dot_operands(pm, True)
+ if amd.has_matrix_core_feature(options.arch):
+ assert options.num_stages != 0, ("Triton AMD backend pipeliner has been updated. "
+ "We used to trigger software pipelining with "
+ "num_stages == 0. Now it will not happen anymore; "
+ "please update to use num_stages == 2 for "
+ "equivalent behavior in the past.")
+ amd.passes.ttgpuir.add_stream_pipelinev2(pm, options.num_stages)
+ passes.common.add_canonicalizer(pm)
+ amd.passes.ttgpuir.insert_instruction_sched_hints(pm)
+ passes.ttgpuir.add_optimize_dot_operands(pm, True)
+ passes.ttgpuir.add_remove_layout_conversions(pm)
+ passes.ttgpuir.add_reduce_data_duplication(pm)
+ if amd.has_matrix_core_feature(options.arch):
+ amd.passes.ttgpuir.add_reorder_instructions(pm)
+ if os.environ.get("AMDGCN_USE_BUFFER_OPS", "0") == "1":
+ amd.passes.ttgpuir.add_canonicalize_pointers(pm)
+ passes.common.add_canonicalizer(pm)
+ amd.passes.ttgpuir.add_convert_to_buffer_ops(pm)
+ passes.common.add_canonicalizer(pm)
+ passes.common.add_cse(pm)
+ passes.common.add_symbol_dce(pm)
+ pm.run(mod)
+ return mod
+
+ @staticmethod
+ def make_llir(src, metadata, options):
+ mod = src
+ # TritonGPU -> LLVM-IR (MLIR)
+ pm = ir.pass_manager(mod.context)
+ pm.enable_debug()
+ amd.passes.ttgpuir.add_decompose_unsupported_conversions(pm, options.arch)
+ # custom_lds_size is an experimental parameter that defines amount of LDS available
+ # for one thread block. Measured in bytes.
+ #
+ # If custom_lds_size = 0, pass will consider all LDS is available for one threads block,
+ # LDS size is determined by provided arch name.
+ custom_lds_size = 0
+ amd.passes.ttgpuir.add_optimize_lds_usage(pm, options.arch, custom_lds_size)
+ passes.convert.add_scf_to_cf(pm)
+ passes.convert.add_index_to_llvmir(pm)
+
+ passes.ttgpuir.add_allocate_shared_memory(pm)
+ ## __HIP_FTZ is used to control the denorm flushing behavior of exp2 op as follows:
+ ## 1. If __HIP_FTZ = 1, exp2 flushes denorms in input and output regardless
+ ## of the value of kernel arg `allow_flush_denorm`.
+ ## 2. If __HIP_FTZ = 0, whether exp2 flushes denorms in input and output
+ ## depends on the value of kernel arg `allow_flush_denorm`.
+ ## 3. __HIP_FTZ is default to 1 and not exposed as a kernel argument.
+ ## For now it is used as a controller for developers only.
+ __HIP_FTZ = True
+ amd.passes.ttgpuir.add_to_llvmir(pm, options.arch, __HIP_FTZ)
+ passes.common.add_canonicalizer(pm)
+ passes.common.add_cse(pm)
+
+ passes.convert.add_cf_to_llvmir(pm)
+ passes.convert.add_arith_to_llvmir(pm)
+ passes.common.add_canonicalizer(pm)
+ passes.common.add_cse(pm)
+ passes.common.add_symbol_dce(pm)
+ amd.passes.ttgpuir.lower_instruction_sched_hints(pm, options.instruction_sched_variant)
+ if os.environ.get("TRITON_DISABLE_LINE_INFO", "0") == "0":
+ passes.llvmir.add_di_scope(pm)
+ amd.passes.ttgpuir.add_builtin_func_to_llvmir(pm, __HIP_FTZ)
+ pm.run(mod)
+
+ # LLVM-IR (MLIR) -> LLVM-IR (LLVM)
+ llvm.init_targets()
+ context = llvm.context()
+ llvm_mod = llvm.to_module(mod, context)
+ amd.attach_target_triple(llvm_mod)
+ llvm.attach_datalayout(llvm_mod, amd.TARGET_TRIPLE, options.arch, '')
+
+ # Set various control constants on the LLVM module so that device
+ # libraries can resolve references to them.
+ amd.set_isa_version(llvm_mod, options.arch)
+ amd.set_abi_version(llvm_mod, 400)
+ amd.set_bool_control_constant(llvm_mod, "__oclc_finite_only_opt", False)
+ amd.set_bool_control_constant(llvm_mod, "__oclc_correctly_rounded_sqrt32", True)
+ amd.set_bool_control_constant(llvm_mod, "__oclc_unsafe_math_opt", False)
+ amd.set_bool_control_constant(llvm_mod, "__oclc_wavefrontsize64", options.warp_size == 64)
+
+ # Set kernel attributes first given this may affect later optimizations.
+ fns = [fn for fn in llvm_mod.get_functions() if not fn.is_declaration()]
+ # The public kernel should be kernel 0.
+ fns[0].set_calling_conv(amd.CALLING_CONV_AMDGPU_KERNEL)
+ fns[0].add_fn_attr("amdgpu-flat-work-group-size", f"1,{options.num_warps*options.warp_size}")
+ fns[0].add_fn_attr("amdgpu-waves-per-eu", f"{options.waves_per_eu}")
+ denormal_mode = "preserve-sign" if options.allow_flush_denorm else "ieee"
+ fns[0].add_fn_attr("denormal-fp-math-f32", denormal_mode)
+
+ # Hint the compiler that we'd like the firmware to set the kernel arguments
+ # to user SGPRs so that the kernel does not need to s_load its arguments
+ # from memory.
+ amd.set_all_fn_arg_inreg(fns[0])
+
+ if options.extern_libs:
+ paths = [path for (name, path) in options.extern_libs if amd.need_extern_lib(llvm_mod, name)]
+ llvm.link_extern_libs(llvm_mod, paths)
+
+ llvm.optimize_module(llvm_mod, llvm.OPTIMIZE_O3, options.arch, '', [], options.enable_fp_fusion)
+
+ # Get some metadata
+ metadata["shared"] = src.get_int_attr("triton_gpu.shared")
+
+ amd.cleanup_bitcode_metadata(llvm_mod)
+ return str(llvm_mod)
+
+ @staticmethod
+ def make_amdgcn(src, metadata, options):
+ # Find kernel names (there should only be one)
+ # We get the name at the last possible step to accomodate `triton.compile`
+ # on user-provided LLVM
+ names = re.findall(r"define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)", src)
+ assert len(names) == 1
+ metadata["name"] = names[0]
+ # llvm -> hsaco
+ amdgcn = llvm.translate_to_asm(src, amd.TARGET_TRIPLE, options.arch, '', [], options.enable_fp_fusion, False)
+ if os.environ.get("AMDGCN_ENABLE_DUMP", "0") == "1":
+ print("// -----// AMDGCN Dump //----- //")
+ print(amdgcn)
+ return amdgcn
+
+ @staticmethod
+ def make_hsaco(src, metadata, options):
+ hsaco = amd.assemble_amdgcn(src, options.arch, '')
+
+ rocm_path = HIPBackend.path_to_rocm_lld()
+ with tempfile.NamedTemporaryFile() as tmp_out:
+ with tempfile.NamedTemporaryFile() as tmp_in:
+ with open(tmp_in.name, 'wb') as fd_in:
+ fd_in.write(hsaco)
+ subprocess.check_call([rocm_path, '-flavor', 'gnu', '-shared', tmp_in.name, '-o', tmp_out.name])
+ with open(tmp_out.name, 'rb') as fd_out:
+ ret = fd_out.read()
+ return ret
+
+ def add_stages(self, stages, options):
+ stages["ttir"] = lambda src, metadata: self.make_ttir(src, metadata, options)
+ stages["ttgir"] = lambda src, metadata: self.make_ttgir(src, metadata, options)
+ stages["llir"] = lambda src, metadata: self.make_llir(src, metadata, options)
+ stages["amdgcn"] = lambda src, metadata: self.make_amdgcn(src, metadata, options)
+ stages["hsaco"] = lambda src, metadata: self.make_hsaco(src, metadata, options)
+
+ @functools.lru_cache()
+ def hash(self):
+ version = subprocess.check_output([HIPBackend.path_to_rocm_lld(), "--version"], encoding='utf-8')
+ return f'{version}-{self.target}'
diff --git a/lib/python3.12/site-packages/triton/backends/amd/driver.c b/lib/python3.12/site-packages/triton/backends/amd/driver.c
new file mode 100644
index 0000000000000000000000000000000000000000..62eee09e78c56291bcfb83e3b12755dc9fcf2c28
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/driver.c
@@ -0,0 +1,211 @@
+#define __HIP_PLATFORM_AMD__
+// clang-format off
+// hip_depreated.h needs definitions from hip_runtime.h.
+#include
+#include
+// clang-format on
+#define PY_SSIZE_T_CLEAN
+#include
+#include
+#include
+#include
+
+// The list of paths to search for the HIP runtime library. The caller Python
+// code should substitute the search path placeholder.
+static const char *hipLibSearchPaths[] = {"/*py_libhip_search_path*/"};
+
+// The list of HIP dynamic library symbols and their signature we are interested
+// in this file.
+// |FOR_EACH_ERR_FN| is a macro to process APIs that return hipError_t;
+// |FOR_EACH_STR_FN| is a macro to process APIs that return const char *.
+//
+// HIP 6.0 introduced an updated hipGetDeviceProperties API under a new symbol,
+// hipGetDevicePropertiesR0600. However, the associated hipDeviceProp_t was
+// directly updated with breaking changes to match hipGetDevicePropertiesR0600
+// in the header file. We include the header file from HIP 6.0. So here if we
+// use hipGetDeviceProperties together with hipDeviceProp_t we will use the
+// old API with a new struct definition and mess up the interpretation.
+//
+// This is a known issue: https://github.com/ROCm/ROCm/issues/2728.
+//
+// For now explicitly defer to the old hipDeviceProp_t struct. This should work
+// for both 5.x and 6.x. In the long term we need to switch to use
+// hipGetProcAddress once available:
+// https://github.com/ROCm/clr/commit/0479cdb3dd30ef58718cad44e424bd793c394cc0
+#define HIP_SYMBOL_LIST(FOR_EACH_ERR_FN, FOR_EACH_STR_FN) \
+ FOR_EACH_STR_FN(hipGetErrorString, hipError_t hipError) \
+ FOR_EACH_ERR_FN(hipGetDeviceProperties, hipDeviceProp_tR0000 *prop, \
+ int deviceId) \
+ FOR_EACH_ERR_FN(hipModuleLoadDataEx, hipModule_t *module, const void *image, \
+ unsigned int numOptions, hipJitOption *options, \
+ void **optionValues) \
+ FOR_EACH_ERR_FN(hipModuleGetFunction, hipFunction_t *function, \
+ hipModule_t module, const char *kname) \
+ FOR_EACH_ERR_FN(hipFuncGetAttribute, int *, hipFunction_attribute attr, \
+ hipFunction_t function)
+
+// The HIP symbol table for holding resolved dynamic library symbols.
+struct HIPSymbolTable {
+#define DEFINE_EACH_ERR_FIELD(hipSymbolName, ...) \
+ hipError_t (*hipSymbolName)(__VA_ARGS__);
+#define DEFINE_EACH_STR_FIELD(hipSymbolName, ...) \
+ const char *(*hipSymbolName)(__VA_ARGS__);
+
+ HIP_SYMBOL_LIST(DEFINE_EACH_ERR_FIELD, DEFINE_EACH_STR_FIELD)
+};
+
+static struct HIPSymbolTable hipSymbolTable;
+
+bool initSymbolTable() {
+ // Use the HIP runtime library loaded into the existing process if it exits.
+ void *lib = dlopen("libamdhip64.so", RTLD_NOLOAD);
+ if (lib) {
+ // printf("[triton] chosen loaded libamdhip64.so in the process\n");
+ }
+
+ // Otherwise, go through the list of search paths to dlopen the first HIP
+ // driver library.
+ if (!lib) {
+ int n = sizeof(hipLibSearchPaths) / sizeof(hipLibSearchPaths[0]);
+ for (int i = 0; i < n; ++i) {
+ void *handle = dlopen(hipLibSearchPaths[i], RTLD_LAZY | RTLD_LOCAL);
+ if (handle) {
+ lib = handle;
+ // printf("[triton] chosen %s\n", hipLibSearchPaths[i]);
+ }
+ }
+ }
+ if (!lib) {
+ PyErr_SetString(PyExc_RuntimeError, "cannot open libamdhip64.so");
+ return false;
+ }
+
+ // Resolve all symbols we are interested in.
+ dlerror(); // Clear existing errors
+ const char *error = NULL;
+#define QUERY_EACH_FN(hipSymbolName, ...) \
+ *(void **)&hipSymbolTable.hipSymbolName = dlsym(lib, #hipSymbolName); \
+ error = dlerror(); \
+ if (error) { \
+ PyErr_SetString(PyExc_RuntimeError, \
+ "cannot query " #hipSymbolName " from libamdhip64.so"); \
+ dlclose(lib); \
+ return false; \
+ }
+
+ HIP_SYMBOL_LIST(QUERY_EACH_FN, QUERY_EACH_FN)
+
+ return true;
+}
+
+static inline void gpuAssert(hipError_t code, const char *file, int line) {
+ {
+ if (code != HIP_SUCCESS) {
+ {
+ const char *prefix = "Triton Error [HIP]: ";
+ const char *str = hipSymbolTable.hipGetErrorString(code);
+ char err[1024] = {0};
+ snprintf(err, 1024, "%s Code: %d, Messsage: %s", prefix, code, str);
+ PyGILState_STATE gil_state;
+ gil_state = PyGILState_Ensure();
+ PyErr_SetString(PyExc_RuntimeError, err);
+ PyGILState_Release(gil_state);
+ }
+ }
+ }
+}
+
+#define HIP_CHECK(ans) \
+ { \
+ gpuAssert((ans), __FILE__, __LINE__); \
+ if (PyErr_Occurred()) \
+ return NULL; \
+ }
+
+static PyObject *getDeviceProperties(PyObject *self, PyObject *args) {
+ int device_id;
+ if (!PyArg_ParseTuple(args, "i", &device_id))
+ return NULL;
+
+ hipDeviceProp_tR0000 props;
+ HIP_CHECK(hipSymbolTable.hipGetDeviceProperties(&props, device_id));
+
+ // create a struct to hold device properties
+ return Py_BuildValue(
+ "{s:i, s:i, s:i, s:i, s:i, s:i, s:s, s:i, s:i}", "max_shared_mem",
+ props.sharedMemPerBlock, "max_num_regs", props.regsPerBlock,
+ "multiprocessor_count", props.multiProcessorCount, "sm_clock_rate",
+ props.clockRate, "mem_clock_rate", props.memoryClockRate, "mem_bus_width",
+ props.memoryBusWidth, "arch", props.gcnArchName, "warpSize",
+ props.warpSize, "max_threads_per_sm", props.maxThreadsPerMultiProcessor);
+}
+
+static PyObject *loadBinary(PyObject *self, PyObject *args) {
+ const char *name;
+ const char *data;
+ Py_ssize_t data_size;
+ int shared;
+ int device;
+ if (!PyArg_ParseTuple(args, "ss#ii", &name, &data, &data_size, &shared,
+ &device)) {
+ return NULL;
+ }
+
+ // set HIP options
+ hipJitOption opt[] = {hipJitOptionErrorLogBufferSizeBytes,
+ hipJitOptionErrorLogBuffer,
+ hipJitOptionInfoLogBufferSizeBytes,
+ hipJitOptionInfoLogBuffer, hipJitOptionLogVerbose};
+ const unsigned int errbufsize = 8192;
+ const unsigned int logbufsize = 8192;
+ char _err[errbufsize];
+ char _log[logbufsize];
+ void *optval[] = {(void *)(uintptr_t)errbufsize, (void *)_err,
+ (void *)(uintptr_t)logbufsize, (void *)_log, (void *)1};
+
+ // launch HIP Binary
+ hipModule_t mod;
+ hipFunction_t fun;
+ HIP_CHECK(hipSymbolTable.hipModuleLoadDataEx(&mod, data, 5, opt, optval))
+ HIP_CHECK(hipSymbolTable.hipModuleGetFunction(&fun, mod, name));
+
+ // get allocated registers and spilled registers from the function
+ int n_regs = 0;
+ int n_spills = 0;
+ hipSymbolTable.hipFuncGetAttribute(&n_regs, HIP_FUNC_ATTRIBUTE_NUM_REGS, fun);
+ hipSymbolTable.hipFuncGetAttribute(&n_spills,
+ HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, fun);
+ n_spills /= 4;
+ if (PyErr_Occurred()) {
+ return NULL;
+ }
+ return Py_BuildValue("(KKii)", (uint64_t)mod, (uint64_t)fun, n_regs,
+ n_spills);
+}
+
+static PyMethodDef ModuleMethods[] = {
+ {"load_binary", loadBinary, METH_VARARGS,
+ "Load provided hsaco into HIP driver"},
+ {"get_device_properties", getDeviceProperties, METH_VARARGS,
+ "Get the properties for a given device"},
+ {NULL, NULL, 0, NULL} // sentinel
+};
+
+static struct PyModuleDef ModuleDef = {PyModuleDef_HEAD_INIT, "hip_utils",
+ NULL, // documentation
+ -1, // size
+ ModuleMethods};
+
+PyMODINIT_FUNC PyInit_hip_utils(void) {
+ if (!initSymbolTable()) {
+ return NULL;
+ }
+
+ PyObject *m = PyModule_Create(&ModuleDef);
+ if (m == NULL) {
+ return NULL;
+ }
+ PyModule_AddFunctions(m, ModuleMethods);
+
+ return m;
+}
diff --git a/lib/python3.12/site-packages/triton/backends/amd/driver.py b/lib/python3.12/site-packages/triton/backends/amd/driver.py
new file mode 100644
index 0000000000000000000000000000000000000000..6e1a368bf8cb1b66ccad7ecbd20b61a4febf7df7
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/driver.py
@@ -0,0 +1,512 @@
+import functools
+import os
+import hashlib
+import subprocess
+import tempfile
+from pathlib import Path
+from triton.runtime.build import _build
+from triton.runtime.cache import get_cache_manager
+from triton.backends.compiler import GPUTarget
+from triton.backends.driver import GPUDriver
+
+dirname = os.path.dirname(os.path.realpath(__file__))
+include_dir = [os.path.join(dirname, "include")]
+
+
+def _find_already_mmapped_dylib_on_linux(lib_name):
+ import platform
+ if platform.system() != 'Linux':
+ return None
+
+ # Use dl_iterate_phdr to walk through the list of shared libraries at runtime.
+ # See https://www.man7.org/linux/man-pages/man3/dl_iterate_phdr.3.html for details.
+
+ import ctypes
+ from ctypes import c_char, c_int, c_size_t, c_void_p, c_char_p, POINTER
+
+ class DlPhdrInfo(ctypes.Structure):
+ _fields_ = [
+ ('dlpi_addr', c_void_p),
+ ('dlpi_name', c_char_p),
+ # We don't care about the remaining fields.
+ ]
+
+ # callback_t must use POINTER(c_char) to avoid copying.
+ callback_t = ctypes.CFUNCTYPE(c_int, POINTER(DlPhdrInfo), POINTER(c_size_t), POINTER(c_char))
+
+ # Load libc and get the dl_iterate_phdr symbol.
+ try:
+ dl_iterate_phdr = ctypes.CDLL('libc.so.6').dl_iterate_phdr
+ except:
+ return None
+ # argtypes must use c_char_p to accept create_string_buffer.
+ dl_iterate_phdr.argtypes = [callback_t, c_char_p]
+ dl_iterate_phdr.restype = c_int
+
+ max_path_length = 4096
+ path = ctypes.create_string_buffer(max_path_length + 1)
+
+ # Define callback to get the loaded dylib path.
+ def callback(info, size, data):
+ dlpi_name = info.contents.dlpi_name
+ p = Path(os.fsdecode(dlpi_name))
+ if lib_name in p.name:
+ # Found the dylib; get its path.
+ ctypes.memmove(data, dlpi_name, min(max_path_length, len(dlpi_name)))
+ return 1
+ return 0
+
+ if dl_iterate_phdr(callback_t(callback), path):
+ return os.fsdecode(ctypes.string_at(path))
+ return None
+
+
+@functools.lru_cache()
+def _get_path_to_hip_runtime_dylib():
+ lib_name = "libamdhip64.so"
+
+ # If we are told explicitly what HIP runtime dynamic library to use, obey that.
+ env_libhip_path = os.getenv("TRITON_LIBHIP_PATH")
+ if env_libhip_path:
+ if env_libhip_path.endswith(lib_name) and os.path.exists(env_libhip_path):
+ return env_libhip_path
+ raise RuntimeError(f"TRITON_LIBHIP_PATH '{env_libhip_path}' does not point to a valid {lib_name}")
+
+ # If the shared object is already mmapped to address space, use it.
+ mmapped_path = _find_already_mmapped_dylib_on_linux(lib_name)
+ if mmapped_path:
+ if os.path.exists(mmapped_path):
+ return mmapped_path
+ raise RuntimeError(f"memory mapped '{mmapped_path}' in process does not point to a valid {lib_name}")
+
+ paths = []
+
+ import site
+ # First search the HIP runtime dynamic library packaged with PyTorch. It's very likely
+ # that we run Triton together with PyTorch. This makes sure we use the same dynamic
+ # library to avoid version mismatch.
+ site_packages = site.getsitepackages()
+ user_site = site.getusersitepackages()
+ if site.ENABLE_USER_SITE: # ENABLE_USER_SITE is initialized in getusersitepackages()
+ site_packages = [user_site] + site_packages
+ for path in site_packages:
+ path = os.path.join(path, "torch", "lib", lib_name)
+ if os.path.exists(path):
+ return path
+ paths.append(path)
+
+ # Then try to see if developer provides a HIP runtime dynamic library using LD_LIBARAY_PATH.
+ env_ld_library_path = os.getenv("LD_LIBRARY_PATH")
+ if env_ld_library_path:
+ for d in env_ld_library_path.split(":"):
+ f = os.path.join(d, lib_name)
+ if os.path.exists(f):
+ return f
+ paths.append(f)
+
+ # Afterwards try to search the loader dynamic library resolution paths.
+ libs = subprocess.check_output(["/sbin/ldconfig", "-p"]).decode()
+ # each line looks like the following:
+ # libamdhip64.so.6 (libc6,x86-64) => /opt/rocm-6.0.2/lib/libamdhip64.so.6
+ # libamdhip64.so (libc6,x86-64) => /opt/rocm-6.0.2/lib/libamdhip64.so
+ locs = [line.split()[-1] for line in libs.splitlines() if line.strip().endswith(lib_name)]
+ for loc in locs:
+ if os.path.exists(loc):
+ return loc
+ paths.append(loc)
+
+ # As a last resort, guess if we have it in some common installation path.
+ common_install_path = os.path.join('/opt/rocm/lib/', lib_name)
+ if os.path.exists(common_install_path):
+ return common_install_path
+ paths.append(common_install_path)
+
+ raise RuntimeError(f"cannot locate {lib_name} after attempted paths {paths}")
+
+
+def compile_module_from_src(src, name):
+ key = hashlib.sha256(src.encode("utf-8")).hexdigest()
+ cache = get_cache_manager(key)
+ cache_path = cache.get_file(f"{name}.so")
+ if cache_path is None:
+ with tempfile.TemporaryDirectory() as tmpdir:
+ src_path = os.path.join(tmpdir, "main.c")
+ with open(src_path, "w") as f:
+ f.write(src)
+ so = _build(name, src_path, tmpdir, [], include_dir, [])
+ with open(so, "rb") as f:
+ cache_path = cache.put(f.read(), f"{name}.so", binary=True)
+ import importlib.util
+ spec = importlib.util.spec_from_file_location(name, cache_path)
+ mod = importlib.util.module_from_spec(spec)
+ spec.loader.exec_module(mod)
+ return mod
+
+
+class HIPUtils(object):
+
+ def __new__(cls):
+ if not hasattr(cls, "instance"):
+ cls.instance = super(HIPUtils, cls).__new__(cls)
+ return cls.instance
+
+ def __init__(self):
+ libhip_path = _get_path_to_hip_runtime_dylib()
+ src = Path(os.path.join(dirname, "driver.c")).read_text()
+ # Just do a simple search and replace here instead of templates or format strings.
+ # This way we don't need to escape-quote C code curly brackets and we can replace
+ # exactly once.
+ src = src.replace('/*py_libhip_search_path*/', libhip_path, 1)
+ mod = compile_module_from_src(src, "hip_utils")
+ self.load_binary = mod.load_binary
+ self.get_device_properties = mod.get_device_properties
+
+
+# -------------------- Launcher ----------------------------
+def ty_to_cpp(ty):
+ if ty[0] == '*':
+ return "hipDeviceptr_t"
+ return {
+ "i1": "int32_t",
+ "i8": "int8_t",
+ "i16": "int16_t",
+ "i32": "int32_t",
+ "i64": "int64_t",
+ "u1": "uint32_t",
+ "u8": "uint8_t",
+ "u16": "uint16_t",
+ "u32": "uint32_t",
+ "u64": "uint64_t",
+ "fp16": "float",
+ "bf16": "float",
+ "fp32": "float",
+ "f32": "float",
+ "fp64": "double",
+ }[ty]
+
+
+def make_launcher(constants, signature, ids, warp_size):
+ start_desc = len(signature)
+ #signature = generate_cu_signature(constants, signature, ids)
+ arg_decls = ', '.join(f"{ty_to_cpp(ty)} arg{i}" for i, ty in signature.items())
+
+ def _extracted_type(ty):
+ if ty[0] == '*':
+ return "PyObject*"
+ return {
+ 'i1': 'int32_t',
+ 'i8': 'int8_t',
+ 'i16': 'int16_t',
+ 'i32': 'int32_t',
+ 'i64': 'int64_t',
+ 'u1': 'uint32_t',
+ 'u8': 'uint8_t',
+ 'u16': 'uint16_t',
+ 'u32': 'uint32_t',
+ 'u64': 'uint64_t',
+ 'fp16': 'float',
+ 'bf16': 'float',
+ 'fp32': 'float',
+ 'f32': 'float',
+ 'fp64': 'double',
+ }[ty]
+
+ def format_of(ty):
+ return {
+ "PyObject*": "O",
+ "float": "f",
+ "double": "d",
+ "long": "l",
+ "int8_t": "b",
+ "int16_t": "h",
+ "int32_t": "i",
+ "int64_t": "l",
+ "uint8_t": "B",
+ "uint16_t": "H",
+ "uint32_t": "I",
+ "uint64_t": "K",
+ }[ty]
+
+ args_format = ''.join([format_of(_extracted_type(ty)) for ty in signature.values()])
+ format = "iiiKKOOOO" + args_format
+ args_list = ', ' + ', '.join(f"&_arg{i}" for i, ty in signature.items()) if len(signature) > 0 else ''
+
+ libhip_path = _get_path_to_hip_runtime_dylib()
+
+ # generate glue code
+ params = [i for i in signature.keys() if i not in constants]
+ src = f"""
+#define __HIP_PLATFORM_AMD__
+#include
+#include
+#include
+#include
+#include
+
+// The list of paths to search for the HIP runtime library. The caller Python
+// code should substitute the search path placeholder.
+static const char *hipLibSearchPaths[] = {{"{libhip_path}"}};
+
+// The list of HIP dynamic library symbols and their signature we are interested
+// in this file.
+#define HIP_SYMBOL_LIST(FOR_EACH_ERR_FN, FOR_EACH_STR_FN) \\
+ FOR_EACH_STR_FN(hipGetErrorString, hipError_t hipError) \\
+ FOR_EACH_ERR_FN(hipModuleLaunchKernel, hipFunction_t f, \\
+ unsigned int gridDimX, unsigned int gridDimY, \\
+ unsigned int gridDimZ, unsigned int blockDimX, \\
+ unsigned int blockDimY, unsigned int blockDimZ, \\
+ unsigned int sharedMemBytes, hipStream_t stream, \\
+ void **kernelParams, void **extra) \\
+ FOR_EACH_ERR_FN(hipPointerGetAttribute, void *data, \\
+ hipPointer_attribute attribute, hipDeviceptr_t ptr)
+
+// The HIP symbol table for holding resolved dynamic library symbols.
+struct HIPSymbolTable {{
+#define DEFINE_EACH_ERR_FIELD(hipSymbolName, ...) \\
+ hipError_t (*hipSymbolName)(__VA_ARGS__);
+#define DEFINE_EACH_STR_FIELD(hipSymbolName, ...) \\
+ const char *(*hipSymbolName)(__VA_ARGS__);
+
+ HIP_SYMBOL_LIST(DEFINE_EACH_ERR_FIELD, DEFINE_EACH_STR_FIELD)
+}};
+
+static struct HIPSymbolTable hipSymbolTable;
+
+bool initSymbolTable() {{
+ // Use the HIP runtime library loaded into the existing process if it exits.
+ void *lib = dlopen("libamdhip64.so", RTLD_NOLOAD);
+ if (lib) {{
+ // printf("[triton] chosen loaded libamdhip64.so in the process\\n");
+ }}
+
+ // Otherwise, go through the list of search paths to dlopen the first HIP
+ // driver library.
+ if (!lib) {{
+ int n = sizeof(hipLibSearchPaths) / sizeof(hipLibSearchPaths[0]);
+ for (int i = 0; i < n; ++i) {{
+ void *handle = dlopen(hipLibSearchPaths[i], RTLD_LAZY | RTLD_LOCAL);
+ if (handle) {{
+ lib = handle;
+ // printf("[triton] chosen %s\\n", hipLibSearchPaths[i]);
+ }}
+ }}
+ }}
+ if (!lib) {{
+ PyErr_SetString(PyExc_RuntimeError, "cannot open libamdhip64.so");
+ return false;
+ }}
+
+ // Resolve all symbols we are interested in.
+ dlerror(); // Clear existing errors
+ const char *error = NULL;
+#define QUERY_EACH_FN(hipSymbolName, ...) \\
+ *(void **)&hipSymbolTable.hipSymbolName = dlsym(lib, #hipSymbolName); \\
+ error = dlerror(); \\
+ if (error) {{ \\
+ PyErr_SetString(PyExc_RuntimeError, \\
+ "cannot query " #hipSymbolName " from libamdhip64.so"); \\
+ dlclose(lib); \\
+ return false; \\
+ }}
+
+ HIP_SYMBOL_LIST(QUERY_EACH_FN, QUERY_EACH_FN)
+
+ return true;
+}}
+
+static inline void gpuAssert(hipError_t code, const char *file, int line)
+{{
+ if (code != HIP_SUCCESS)
+ {{
+ const char* prefix = "Triton Error [HIP]: ";
+ const char* str = hipSymbolTable.hipGetErrorString(code);
+ char err[1024] = {{0}};
+ snprintf(err, 1024, "%s Code: %d, Messsage: %s", prefix, code, str );
+ PyErr_SetString(PyExc_RuntimeError, err);
+ }}
+}}
+
+#define HIP_CHECK(ans) {{ gpuAssert((ans), __FILE__, __LINE__); }}
+
+static void _launch(int gridX, int gridY, int gridZ, int num_warps, int num_ctas, int clusterDimX, int clusterDimY, int clusterDimZ, int shared_memory, hipStream_t stream, hipFunction_t function{', ' + arg_decls if len(arg_decls) > 0 else ''}) {{
+ // printf("_launch hip kernel\\n");
+ void *params[] = {{ {', '.join(f"&arg{i}" for i in params)} }};
+ if (gridX*gridY*gridZ > 0) {{
+ HIP_CHECK(hipSymbolTable.hipModuleLaunchKernel(function, gridX, gridY, gridZ, {warp_size}*num_warps, 1, 1, shared_memory, stream, params, 0));
+ }}
+ }}
+
+typedef struct _DevicePtrInfo {{
+ hipDeviceptr_t dev_ptr;
+ bool valid;
+}} DevicePtrInfo;
+
+static inline DevicePtrInfo getPointer(PyObject *obj, int idx) {{
+ DevicePtrInfo ptr_info;
+ ptr_info.dev_ptr = 0;
+ ptr_info.valid = true;
+ if (PyLong_Check(obj)) {{
+ ptr_info.dev_ptr = (hipDeviceptr_t)PyLong_AsUnsignedLongLong(obj);
+ return ptr_info;
+ }}
+ if (obj == Py_None) {{
+ // valid nullptr
+ return ptr_info;
+ }}
+ PyObject *ptr = PyObject_GetAttrString(obj, "data_ptr");
+ if(ptr){{
+ PyObject *empty_tuple = PyTuple_New(0);
+ PyObject *ret = PyObject_Call(ptr, empty_tuple, NULL);
+ Py_DECREF(empty_tuple);
+ Py_DECREF(ptr);
+ if (!PyLong_Check(ret)) {{
+ PyErr_SetString(PyExc_TypeError, "data_ptr method of Pointer object must return 64-bit int");
+ ptr_info.valid = false;
+ return ptr_info;
+ }}
+ ptr_info.dev_ptr = (hipDeviceptr_t)PyLong_AsUnsignedLongLong(ret);
+ if(!ptr_info.dev_ptr)
+ return ptr_info;
+ uint64_t dev_ptr;
+ hipError_t status = hipSymbolTable.hipPointerGetAttribute(&dev_ptr, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, ptr_info.dev_ptr);
+ if (status == hipErrorInvalidValue) {{
+ PyErr_Format(PyExc_ValueError,
+ "Pointer argument (at %d) cannot be accessed from Triton (cpu tensor?)", idx);
+ ptr_info.valid = false;
+ }}
+ ptr_info.dev_ptr = (hipDeviceptr_t)dev_ptr;
+ Py_DECREF(ret);
+ return ptr_info;
+ }}
+ PyErr_SetString(PyExc_TypeError, "Pointer argument must be either uint64 or have data_ptr method");
+ return ptr_info;
+}}
+
+static PyObject* launch(PyObject* self, PyObject* args) {{
+ // printf("launch\\n");
+ int gridX, gridY, gridZ;
+ uint64_t _stream;
+ uint64_t _function;
+ PyObject *launch_enter_hook = NULL;
+ PyObject *launch_exit_hook = NULL;
+ PyObject *kernel_metadata = NULL;
+ PyObject *launch_metadata = NULL;
+ {' '.join([f"{_extracted_type(ty)} _arg{i}; " for i, ty in signature.items()])}
+ if(!PyArg_ParseTuple(args, \"{format}\", &gridX, &gridY, &gridZ, &_stream, &_function,
+ &kernel_metadata, &launch_metadata,
+ &launch_enter_hook, &launch_exit_hook {args_list})) {{
+ return NULL;
+ }}
+
+ // extract kernel metadata
+ int num_warps, num_ctas, shared_memory, clusterDimX, clusterDimY, clusterDimZ;
+ if (!PyArg_ParseTuple(kernel_metadata, \"iiiiii\", &num_warps, &num_ctas, &shared_memory, &clusterDimX, &clusterDimY, &clusterDimZ)) {{
+ return NULL;
+ }}
+ // extract launch metadata
+ if (launch_enter_hook != Py_None){{
+ PyObject* args = Py_BuildValue("(O)", launch_metadata);
+ PyObject* ret = PyObject_CallObject(launch_enter_hook, args);
+ Py_DECREF(args);
+ if (!ret)
+ return NULL;
+ }}
+
+
+ // raise exception asap
+ {"; ".join([f"DevicePtrInfo ptr_info{i} = getPointer(_arg{i}, {i}); if (!ptr_info{i}.valid) return NULL;" if ty[0] == "*" else "" for i, ty in signature.items()])};
+ _launch(gridX, gridY, gridZ, num_warps, num_ctas, clusterDimX, clusterDimY, clusterDimZ, shared_memory, (hipStream_t)_stream, (hipFunction_t)_function{', ' + ', '.join(f"ptr_info{i}.dev_ptr" if ty[0]=="*" else f"_arg{i}"for i, ty in signature.items()) if len(signature) > 0 else ''});
+
+ if(launch_exit_hook != Py_None){{
+ PyObject* args = Py_BuildValue("(O)", launch_metadata);
+ PyObject* ret = PyObject_CallObject(launch_exit_hook, args);
+ Py_DECREF(args);
+ if (!ret)
+ return NULL;
+ }}
+
+ if(PyErr_Occurred()) {{
+ return NULL;
+ }}
+ // return None
+ Py_INCREF(Py_None);
+ return Py_None;
+}}
+
+static PyMethodDef ModuleMethods[] = {{
+ {{"launch", launch, METH_VARARGS, "Entry point for all kernels with this signature"}},
+ {{NULL, NULL, 0, NULL}} // sentinel
+}};
+
+static struct PyModuleDef ModuleDef = {{
+ PyModuleDef_HEAD_INIT,
+ \"__triton_launcher\",
+ NULL, //documentation
+ -1, //size
+ ModuleMethods
+}};
+
+PyMODINIT_FUNC PyInit___triton_launcher(void) {{
+ if (!initSymbolTable()) {{
+ return NULL;
+ }}
+ PyObject *m = PyModule_Create(&ModuleDef);
+ if(m == NULL) {{
+ return NULL;
+ }}
+ PyModule_AddFunctions(m, ModuleMethods);
+ return m;
+}}
+"""
+ return src
+
+
+class HIPLauncher(object):
+
+ def __init__(self, src, metadata):
+ ids = {"ids_of_const_exprs": src.fn.constexprs if hasattr(src, "fn") else tuple()}
+ constants = src.constants if hasattr(src, "constants") else dict()
+ cst_key = lambda i: src.fn.arg_names.index(i) if isinstance(i, str) else i
+ constants = {cst_key(key): value for key, value in constants.items()}
+ signature = {cst_key(key): value for key, value in src.signature.items()}
+ src = make_launcher(constants, signature, ids, metadata.warp_size)
+ mod = compile_module_from_src(src, "__triton_launcher")
+ self.launch = mod.launch
+
+ def __call__(self, *args, **kwargs):
+ self.launch(*args, **kwargs)
+
+
+class HIPDriver(GPUDriver):
+
+ def __init__(self):
+ super().__init__()
+ self.utils = HIPUtils()
+ self.launcher_cls = HIPLauncher
+
+ def get_device_interface(self):
+ import torch
+ return torch.cuda
+
+ @staticmethod
+ def is_active():
+ import torch
+ return torch.version.hip is not None
+
+ def get_current_target(self):
+ device = self.get_current_device()
+ device_properties = self.utils.get_device_properties(device)
+ arch = device_properties['arch']
+ warp_size = device_properties['warpSize']
+ return GPUTarget("hip", arch.split(':')[0], warp_size)
+
+ def get_benchmarker(self):
+ from triton.testing import do_bench
+ return do_bench
+
+ def get_empty_cache_for_benchmark(self):
+ import torch
+
+ # It's the same as the Nvidia backend.
+ cache_size = 256 * 1024 * 1024
+ return torch.empty(int(cache_size // 4), dtype=torch.int, device='cuda')
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_channel_descriptor.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_channel_descriptor.h
new file mode 100644
index 0000000000000000000000000000000000000000..c6b150d4b32237468a947b9c0e31843560751357
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_channel_descriptor.h
@@ -0,0 +1,358 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
+#define HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H
+
+#if !defined(__HIPCC_RTC__)
+#include
+#include
+#include
+#endif
+
+#ifdef __cplusplus
+
+extern "C" HIP_PUBLIC_API
+hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
+
+static inline hipChannelFormatDesc hipCreateChannelDescHalf() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
+}
+
+static inline hipChannelFormatDesc hipCreateChannelDescHalf1() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
+}
+
+static inline hipChannelFormatDesc hipCreateChannelDescHalf2() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindFloat);
+}
+
+static inline hipChannelFormatDesc hipCreateChannelDescHalf4() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindFloat);
+}
+
+template
+static inline hipChannelFormatDesc hipCreateChannelDesc() {
+ return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(char) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed char) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned char) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned char) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed char) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned char) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed char) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
+}
+
+#ifndef __GNUC__ // vector3 is the same as vector4
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned char) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed char) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
+}
+#endif
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned char) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed char) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed short) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed short) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed short) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
+}
+
+#ifndef __GNUC__
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed short) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
+}
+#endif
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned short) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed short) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned int) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed int) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned int) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed int) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned int) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed int) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
+}
+
+#ifndef __GNUC__
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned int) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed int) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
+}
+#endif
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned int) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed int) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(float) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(float) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(float) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindFloat);
+}
+
+#ifndef __GNUC__
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(float) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindFloat);
+}
+#endif
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(float) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindFloat);
+}
+
+#if !defined(__LP64__)
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned long) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed long) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned long) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed long) * 8;
+ return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned long) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed long) * 8;
+ return hipCreateChannelDesc(e, e, 0, 0, hipChannelFormatKindSigned);
+}
+
+#ifndef __GNUC__
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned long) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed long) * 8;
+ return hipCreateChannelDesc(e, e, e, 0, hipChannelFormatKindSigned);
+}
+#endif
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(unsigned long) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindUnsigned);
+}
+
+template <>
+inline hipChannelFormatDesc hipCreateChannelDesc() {
+ int e = (int)sizeof(signed long) * 8;
+ return hipCreateChannelDesc(e, e, e, e, hipChannelFormatKindSigned);
+}
+#endif /* !__LP64__ */
+
+#else
+
+struct hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w,
+ enum hipChannelFormatKind f);
+
+#endif /* __cplusplus */
+
+#endif /* !HIP_INCLUDE_HIP_AMD_DETAIL_CHANNEL_DESCRIPTOR_H */
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_device_functions.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_device_functions.h
new file mode 100644
index 0000000000000000000000000000000000000000..2a5cf48a039771144d87b1ba6d8395e52a5ac90e
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_device_functions.h
@@ -0,0 +1,1031 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_FUNCTIONS_H
+#define HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_FUNCTIONS_H
+
+#if !defined(__HIPCC_RTC__)
+#include
+#include
+#include
+#include "host_defines.h"
+#include "math_fwd.h"
+#include
+#include
+#include
+#endif // !defined(__HIPCC_RTC__)
+
+#if defined(__clang__) && defined(__HIP__)
+extern "C" __device__ int printf(const char *fmt, ...);
+#else
+template
+static inline __device__ void printf(const char* format, All... all) {}
+#endif
+
+extern "C" __device__ unsigned long long __ockl_steadyctr_u64();
+
+/*
+Integer Intrinsics
+*/
+
+// integer intrinsic function __poc __clz __ffs __brev
+__device__ static inline unsigned int __popc(unsigned int input) {
+ return __builtin_popcount(input);
+}
+__device__ static inline unsigned int __popcll(unsigned long long int input) {
+ return __builtin_popcountll(input);
+}
+
+__device__ static inline int __clz(int input) {
+ return __ockl_clz_u32((uint)input);
+}
+
+__device__ static inline int __clzll(long long int input) {
+ return __ockl_clz_u64((uint64_t)input);
+}
+
+__device__ static inline unsigned int __ffs(unsigned int input) {
+ return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
+}
+
+__device__ static inline unsigned int __ffsll(unsigned long long int input) {
+ return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
+}
+
+__device__ static inline unsigned int __ffs(int input) {
+ return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
+}
+
+__device__ static inline unsigned int __ffsll(long long int input) {
+ return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
+}
+
+// Given a 32/64-bit value exec mask and an integer value base (between 0 and WAVEFRONT_SIZE),
+// find the n-th (given by offset) set bit in the exec mask from the base bit, and return the bit position.
+// If not found, return -1.
+__device__ static int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset) {
+ uint64_t temp_mask = mask;
+ int32_t temp_offset = offset;
+
+ if (offset == 0) {
+ temp_mask &= (1 << base);
+ temp_offset = 1;
+ }
+ else if (offset < 0) {
+ temp_mask = __builtin_bitreverse64(mask);
+ base = 63 - base;
+ temp_offset = -offset;
+ }
+
+ temp_mask = temp_mask & ((~0ULL) << base);
+ if (__builtin_popcountll(temp_mask) < temp_offset)
+ return -1;
+ int32_t total = 0;
+ for (int i = 0x20; i > 0; i >>= 1) {
+ uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
+ int32_t pcnt = __builtin_popcountll(temp_mask_lo);
+ if (pcnt < temp_offset) {
+ temp_mask = temp_mask >> i;
+ temp_offset -= pcnt;
+ total += i;
+ }
+ else {
+ temp_mask = temp_mask_lo;
+ }
+ }
+ if (offset < 0)
+ return 63 - total;
+ else
+ return total;
+}
+
+__device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset) {
+ uint64_t temp_mask = mask;
+ int32_t temp_offset = offset;
+ if (offset == 0) {
+ temp_mask &= (1 << base);
+ temp_offset = 1;
+ }
+ else if (offset < 0) {
+ temp_mask = __builtin_bitreverse64(mask);
+ base = 63 - base;
+ temp_offset = -offset;
+ }
+ temp_mask = temp_mask & ((~0ULL) << base);
+ if (__builtin_popcountll(temp_mask) < temp_offset)
+ return -1;
+ int32_t total = 0;
+ for (int i = 0x20; i > 0; i >>= 1) {
+ uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
+ int32_t pcnt = __builtin_popcountll(temp_mask_lo);
+ if (pcnt < temp_offset) {
+ temp_mask = temp_mask >> i;
+ temp_offset -= pcnt;
+ total += i;
+ }
+ else {
+ temp_mask = temp_mask_lo;
+ }
+ }
+ if (offset < 0)
+ return 63 - total;
+ else
+ return total;
+}
+__device__ static inline unsigned int __brev(unsigned int input) {
+ return __builtin_bitreverse32(input);
+}
+
+__device__ static inline unsigned long long int __brevll(unsigned long long int input) {
+ return __builtin_bitreverse64(input);
+}
+
+__device__ static inline unsigned int __lastbit_u32_u64(uint64_t input) {
+ return input == 0 ? -1 : __builtin_ctzl(input);
+}
+
+__device__ static inline unsigned int __bitextract_u32(unsigned int src0, unsigned int src1, unsigned int src2) {
+ uint32_t offset = src1 & 31;
+ uint32_t width = src2 & 31;
+ return width == 0 ? 0 : (src0 << (32 - offset - width)) >> (32 - width);
+}
+
+__device__ static inline uint64_t __bitextract_u64(uint64_t src0, unsigned int src1, unsigned int src2) {
+ uint64_t offset = src1 & 63;
+ uint64_t width = src2 & 63;
+ return width == 0 ? 0 : (src0 << (64 - offset - width)) >> (64 - width);
+}
+
+__device__ static inline unsigned int __bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int src2, unsigned int src3) {
+ uint32_t offset = src2 & 31;
+ uint32_t width = src3 & 31;
+ uint32_t mask = (1 << width) - 1;
+ return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
+}
+
+__device__ static inline uint64_t __bitinsert_u64(uint64_t src0, uint64_t src1, unsigned int src2, unsigned int src3) {
+ uint64_t offset = src2 & 63;
+ uint64_t width = src3 & 63;
+ uint64_t mask = (1ULL << width) - 1;
+ return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
+}
+
+__device__ inline unsigned int __funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift)
+{
+ uint32_t mask_shift = shift & 31;
+ return mask_shift == 0 ? hi : __builtin_amdgcn_alignbit(hi, lo, 32 - mask_shift);
+}
+
+__device__ inline unsigned int __funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift)
+{
+ uint32_t min_shift = shift >= 32 ? 32 : shift;
+ return min_shift == 0 ? hi : __builtin_amdgcn_alignbit(hi, lo, 32 - min_shift);
+}
+
+__device__ inline unsigned int __funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift)
+{
+ return __builtin_amdgcn_alignbit(hi, lo, shift);
+}
+
+__device__ inline unsigned int __funnelshift_rc(unsigned int lo, unsigned int hi, unsigned int shift)
+{
+ return shift >= 32 ? hi : __builtin_amdgcn_alignbit(hi, lo, shift);
+}
+
+__device__ static unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s);
+__device__ static unsigned int __hadd(int x, int y);
+__device__ static int __mul24(int x, int y);
+__device__ static long long int __mul64hi(long long int x, long long int y);
+__device__ static int __mulhi(int x, int y);
+__device__ static int __rhadd(int x, int y);
+__device__ static unsigned int __sad(int x, int y,unsigned int z);
+__device__ static unsigned int __uhadd(unsigned int x, unsigned int y);
+__device__ static int __umul24(unsigned int x, unsigned int y);
+__device__ static unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y);
+__device__ static unsigned int __umulhi(unsigned int x, unsigned int y);
+__device__ static unsigned int __urhadd(unsigned int x, unsigned int y);
+__device__ static unsigned int __usad(unsigned int x, unsigned int y, unsigned int z);
+
+struct ucharHolder {
+ union {
+ unsigned char c[4];
+ unsigned int ui;
+ };
+} __attribute__((aligned(4)));
+
+struct uchar2Holder {
+ union {
+ unsigned int ui[2];
+ unsigned char c[8];
+ };
+} __attribute__((aligned(8)));
+
+__device__
+static inline unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) {
+ struct uchar2Holder cHoldVal;
+ struct ucharHolder cHoldKey;
+ cHoldKey.ui = s;
+ cHoldVal.ui[0] = x;
+ cHoldVal.ui[1] = y;
+ unsigned int result;
+ result = cHoldVal.c[cHoldKey.c[0] & 0x07];
+ result += (cHoldVal.c[(cHoldKey.c[0] & 0x70) >> 4] << 8);
+ result += (cHoldVal.c[cHoldKey.c[1] & 0x07] << 16);
+ result += (cHoldVal.c[(cHoldKey.c[1] & 0x70) >> 4] << 24);
+ return result;
+}
+
+__device__ static inline unsigned int __hadd(int x, int y) {
+ int z = x + y;
+ int sign = z & 0x8000000;
+ int value = z & 0x7FFFFFFF;
+ return ((value) >> 1 || sign);
+}
+
+__device__ static inline int __mul24(int x, int y) {
+ return __ockl_mul24_i32(x, y);
+}
+
+__device__ static inline long long __mul64hi(long long int x, long long int y) {
+ ulong x0 = (ulong)x & 0xffffffffUL;
+ long x1 = x >> 32;
+ ulong y0 = (ulong)y & 0xffffffffUL;
+ long y1 = y >> 32;
+ ulong z0 = x0*y0;
+ long t = x1*y0 + (z0 >> 32);
+ long z1 = t & 0xffffffffL;
+ long z2 = t >> 32;
+ z1 = x0*y1 + z1;
+ return x1*y1 + z2 + (z1 >> 32);
+}
+
+__device__ static inline int __mulhi(int x, int y) {
+ return __ockl_mul_hi_i32(x, y);
+}
+
+__device__ static inline int __rhadd(int x, int y) {
+ int z = x + y + 1;
+ int sign = z & 0x8000000;
+ int value = z & 0x7FFFFFFF;
+ return ((value) >> 1 || sign);
+}
+__device__ static inline unsigned int __sad(int x, int y, unsigned int z) {
+ return x > y ? x - y + z : y - x + z;
+}
+__device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) {
+ return (x + y) >> 1;
+}
+__device__ static inline int __umul24(unsigned int x, unsigned int y) {
+ return __ockl_mul24_u32(x, y);
+}
+
+__device__
+static inline unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) {
+ ulong x0 = x & 0xffffffffUL;
+ ulong x1 = x >> 32;
+ ulong y0 = y & 0xffffffffUL;
+ ulong y1 = y >> 32;
+ ulong z0 = x0*y0;
+ ulong t = x1*y0 + (z0 >> 32);
+ ulong z1 = t & 0xffffffffUL;
+ ulong z2 = t >> 32;
+ z1 = x0*y1 + z1;
+ return x1*y1 + z2 + (z1 >> 32);
+}
+
+__device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y) {
+ return __ockl_mul_hi_u32(x, y);
+}
+__device__ static inline unsigned int __urhadd(unsigned int x, unsigned int y) {
+ return (x + y + 1) >> 1;
+}
+__device__ static inline unsigned int __usad(unsigned int x, unsigned int y, unsigned int z) {
+ return __ockl_sadd_u32(x, y, z);
+}
+
+__device__ static inline unsigned int __lane_id() {
+ return __builtin_amdgcn_mbcnt_hi(
+ -1, __builtin_amdgcn_mbcnt_lo(-1, 0));
+}
+
+__device__
+static inline unsigned int __mbcnt_lo(unsigned int x, unsigned int y) {return __builtin_amdgcn_mbcnt_lo(x,y);};
+
+__device__
+static inline unsigned int __mbcnt_hi(unsigned int x, unsigned int y) {return __builtin_amdgcn_mbcnt_hi(x,y);};
+
+/*
+HIP specific device functions
+*/
+
+#if !defined(__HIPCC_RTC__)
+#include "amd_warp_functions.h"
+#endif
+
+#define MASK1 0x00ff00ff
+#define MASK2 0xff00ff00
+
+__device__ static inline char4 __hip_hc_add8pk(char4 in1, char4 in2) {
+ char4 out;
+ unsigned one1 = in1.w & MASK1;
+ unsigned one2 = in2.w & MASK1;
+ out.w = (one1 + one2) & MASK1;
+ one1 = in1.w & MASK2;
+ one2 = in2.w & MASK2;
+ out.w = out.w | ((one1 + one2) & MASK2);
+ return out;
+}
+
+__device__ static inline char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
+ char4 out;
+ unsigned one1 = in1.w & MASK1;
+ unsigned one2 = in2.w & MASK1;
+ out.w = (one1 - one2) & MASK1;
+ one1 = in1.w & MASK2;
+ one2 = in2.w & MASK2;
+ out.w = out.w | ((one1 - one2) & MASK2);
+ return out;
+}
+
+__device__ static inline char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
+ char4 out;
+ unsigned one1 = in1.w & MASK1;
+ unsigned one2 = in2.w & MASK1;
+ out.w = (one1 * one2) & MASK1;
+ one1 = in1.w & MASK2;
+ one2 = in2.w & MASK2;
+ out.w = out.w | ((one1 * one2) & MASK2);
+ return out;
+}
+
+__device__ static inline float __double2float_rd(double x) {
+ return __ocml_cvtrtn_f32_f64(x);
+}
+__device__ static inline float __double2float_rn(double x) { return x; }
+__device__ static inline float __double2float_ru(double x) {
+ return __ocml_cvtrtp_f32_f64(x);
+}
+__device__ static inline float __double2float_rz(double x) {
+ return __ocml_cvtrtz_f32_f64(x);
+}
+
+__device__ static inline int __double2hiint(double x) {
+ static_assert(sizeof(double) == 2 * sizeof(int), "");
+
+ int tmp[2];
+ __builtin_memcpy(tmp, &x, sizeof(tmp));
+
+ return tmp[1];
+}
+__device__ static inline int __double2loint(double x) {
+ static_assert(sizeof(double) == 2 * sizeof(int), "");
+
+ int tmp[2];
+ __builtin_memcpy(tmp, &x, sizeof(tmp));
+
+ return tmp[0];
+}
+
+__device__ static inline int __double2int_rd(double x) { return (int)__ocml_floor_f64(x); }
+__device__ static inline int __double2int_rn(double x) { return (int)__ocml_rint_f64(x); }
+__device__ static inline int __double2int_ru(double x) { return (int)__ocml_ceil_f64(x); }
+__device__ static inline int __double2int_rz(double x) { return (int)x; }
+
+__device__ static inline long long int __double2ll_rd(double x) {
+ return (long long)__ocml_floor_f64(x);
+}
+__device__ static inline long long int __double2ll_rn(double x) {
+ return (long long)__ocml_rint_f64(x);
+}
+__device__ static inline long long int __double2ll_ru(double x) {
+ return (long long)__ocml_ceil_f64(x);
+}
+__device__ static inline long long int __double2ll_rz(double x) { return (long long)x; }
+
+__device__ static inline unsigned int __double2uint_rd(double x) {
+ return (unsigned int)__ocml_floor_f64(x);
+}
+__device__ static inline unsigned int __double2uint_rn(double x) {
+ return (unsigned int)__ocml_rint_f64(x);
+}
+__device__ static inline unsigned int __double2uint_ru(double x) {
+ return (unsigned int)__ocml_ceil_f64(x);
+}
+__device__ static inline unsigned int __double2uint_rz(double x) { return (unsigned int)x; }
+
+__device__ static inline unsigned long long int __double2ull_rd(double x) {
+ return (unsigned long long int)__ocml_floor_f64(x);
+}
+__device__ static inline unsigned long long int __double2ull_rn(double x) {
+ return (unsigned long long int)__ocml_rint_f64(x);
+}
+__device__ static inline unsigned long long int __double2ull_ru(double x) {
+ return (unsigned long long int)__ocml_ceil_f64(x);
+}
+__device__ static inline unsigned long long int __double2ull_rz(double x) {
+ return (unsigned long long int)x;
+}
+__device__ static inline long long int __double_as_longlong(double x) {
+ static_assert(sizeof(long long) == sizeof(double), "");
+
+ long long tmp;
+ __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+ return tmp;
+}
+
+/*
+__device__ unsigned short __float2half_rn(float x);
+__device__ float __half2float(unsigned short);
+
+The above device function are not a valid .
+Use
+__device__ __half __float2half_rn(float x);
+__device__ float __half2float(__half);
+from hip_fp16.h
+
+CUDA implements half as unsigned short whereas, HIP doesn't.
+
+*/
+
+__device__ static inline int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); }
+__device__ static inline int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); }
+__device__ static inline int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); }
+__device__ static inline int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); }
+
+__device__ static inline long long int __float2ll_rd(float x) {
+ return (long long int)__ocml_floor_f32(x);
+}
+__device__ static inline long long int __float2ll_rn(float x) {
+ return (long long int)__ocml_rint_f32(x);
+}
+__device__ static inline long long int __float2ll_ru(float x) {
+ return (long long int)__ocml_ceil_f32(x);
+}
+__device__ static inline long long int __float2ll_rz(float x) { return (long long int)x; }
+
+__device__ static inline unsigned int __float2uint_rd(float x) {
+ return (unsigned int)__ocml_floor_f32(x);
+}
+__device__ static inline unsigned int __float2uint_rn(float x) {
+ return (unsigned int)__ocml_rint_f32(x);
+}
+__device__ static inline unsigned int __float2uint_ru(float x) {
+ return (unsigned int)__ocml_ceil_f32(x);
+}
+__device__ static inline unsigned int __float2uint_rz(float x) { return (unsigned int)x; }
+
+__device__ static inline unsigned long long int __float2ull_rd(float x) {
+ return (unsigned long long int)__ocml_floor_f32(x);
+}
+__device__ static inline unsigned long long int __float2ull_rn(float x) {
+ return (unsigned long long int)__ocml_rint_f32(x);
+}
+__device__ static inline unsigned long long int __float2ull_ru(float x) {
+ return (unsigned long long int)__ocml_ceil_f32(x);
+}
+__device__ static inline unsigned long long int __float2ull_rz(float x) {
+ return (unsigned long long int)x;
+}
+
+__device__ static inline int __float_as_int(float x) {
+ static_assert(sizeof(int) == sizeof(float), "");
+
+ int tmp;
+ __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+ return tmp;
+}
+
+__device__ static inline unsigned int __float_as_uint(float x) {
+ static_assert(sizeof(unsigned int) == sizeof(float), "");
+
+ unsigned int tmp;
+ __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+ return tmp;
+}
+
+__device__ static inline double __hiloint2double(int hi, int lo) {
+ static_assert(sizeof(double) == sizeof(uint64_t), "");
+
+ uint64_t tmp0 = (static_cast(hi) << 32ull) | static_cast(lo);
+ double tmp1;
+ __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0));
+
+ return tmp1;
+}
+
+__device__ static inline double __int2double_rn(int x) { return (double)x; }
+
+__device__ static inline float __int2float_rd(int x) {
+ return __ocml_cvtrtn_f32_s32(x);
+}
+__device__ static inline float __int2float_rn(int x) { return (float)x; }
+__device__ static inline float __int2float_ru(int x) {
+ return __ocml_cvtrtp_f32_s32(x);
+}
+__device__ static inline float __int2float_rz(int x) {
+ return __ocml_cvtrtz_f32_s32(x);
+}
+
+__device__ static inline float __int_as_float(int x) {
+ static_assert(sizeof(float) == sizeof(int), "");
+
+ float tmp;
+ __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+ return tmp;
+}
+
+__device__ static inline double __ll2double_rd(long long int x) {
+ return __ocml_cvtrtn_f64_s64(x);
+}
+__device__ static inline double __ll2double_rn(long long int x) { return (double)x; }
+__device__ static inline double __ll2double_ru(long long int x) {
+ return __ocml_cvtrtp_f64_s64(x);
+}
+__device__ static inline double __ll2double_rz(long long int x) {
+ return __ocml_cvtrtz_f64_s64(x);
+}
+
+__device__ static inline float __ll2float_rd(long long int x) {
+ return __ocml_cvtrtn_f32_s64(x);
+}
+__device__ static inline float __ll2float_rn(long long int x) { return (float)x; }
+__device__ static inline float __ll2float_ru(long long int x) {
+ return __ocml_cvtrtp_f32_s64(x);
+}
+__device__ static inline float __ll2float_rz(long long int x) {
+ return __ocml_cvtrtz_f32_s64(x);
+}
+
+__device__ static inline double __longlong_as_double(long long int x) {
+ static_assert(sizeof(double) == sizeof(long long), "");
+
+ double tmp;
+ __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+ return tmp;
+}
+
+__device__ static inline double __uint2double_rn(unsigned int x) { return (double)x; }
+
+__device__ static inline float __uint2float_rd(unsigned int x) {
+ return __ocml_cvtrtn_f32_u32(x);
+}
+__device__ static inline float __uint2float_rn(unsigned int x) { return (float)x; }
+__device__ static inline float __uint2float_ru(unsigned int x) {
+ return __ocml_cvtrtp_f32_u32(x);
+}
+__device__ static inline float __uint2float_rz(unsigned int x) {
+ return __ocml_cvtrtz_f32_u32(x);
+}
+
+__device__ static inline float __uint_as_float(unsigned int x) {
+ static_assert(sizeof(float) == sizeof(unsigned int), "");
+
+ float tmp;
+ __builtin_memcpy(&tmp, &x, sizeof(tmp));
+
+ return tmp;
+}
+
+__device__ static inline double __ull2double_rd(unsigned long long int x) {
+ return __ocml_cvtrtn_f64_u64(x);
+}
+__device__ static inline double __ull2double_rn(unsigned long long int x) { return (double)x; }
+__device__ static inline double __ull2double_ru(unsigned long long int x) {
+ return __ocml_cvtrtp_f64_u64(x);
+}
+__device__ static inline double __ull2double_rz(unsigned long long int x) {
+ return __ocml_cvtrtz_f64_u64(x);
+}
+
+__device__ static inline float __ull2float_rd(unsigned long long int x) {
+ return __ocml_cvtrtn_f32_u64(x);
+}
+__device__ static inline float __ull2float_rn(unsigned long long int x) { return (float)x; }
+__device__ static inline float __ull2float_ru(unsigned long long int x) {
+ return __ocml_cvtrtp_f32_u64(x);
+}
+__device__ static inline float __ull2float_rz(unsigned long long int x) {
+ return __ocml_cvtrtz_f32_u64(x);
+}
+
+#if defined(__clang__) && defined(__HIP__)
+
+// Clock functions
+__device__ long long int __clock64();
+__device__ long long int __clock();
+__device__ long long int clock64();
+__device__ long long int clock();
+__device__ long long int wall_clock64();
+// hip.amdgcn.bc - named sync
+__device__ void __named_sync();
+
+#ifdef __HIP_DEVICE_COMPILE__
+
+// Clock function to return GPU core cycle count.
+// GPU can change its core clock frequency at runtime. The maximum frequency can be queried
+// through hipDeviceAttributeClockRate attribute.
+__device__
+inline __attribute((always_inline))
+long long int __clock64() {
+#if __has_builtin(__builtin_amdgcn_s_memtime)
+ // Exists on gfx8, gfx9, gfx10.1, gfx10.2, gfx10.3
+ return (long long int) __builtin_amdgcn_s_memtime();
+#else
+ // Subject to change when better solution available
+ return (long long int) __builtin_readcyclecounter();
+#endif
+}
+
+__device__
+inline __attribute((always_inline))
+long long int __clock() { return __clock64(); }
+
+// Clock function to return wall clock count at a constant frequency that can be queried
+// through hipDeviceAttributeWallClockRate attribute.
+__device__
+inline __attribute__((always_inline))
+long long int wall_clock64() {
+ return (long long int) __ockl_steadyctr_u64();
+}
+
+__device__
+inline __attribute__((always_inline))
+long long int clock64() { return __clock64(); }
+
+__device__
+inline __attribute__((always_inline))
+long long int clock() { return __clock(); }
+
+// hip.amdgcn.bc - named sync
+__device__
+inline
+void __named_sync() { __builtin_amdgcn_s_barrier(); }
+
+#endif // __HIP_DEVICE_COMPILE__
+
+// warp vote function __all __any __ballot
+__device__
+inline
+int __all(int predicate) {
+ return __ockl_wfall_i32(predicate);
+}
+
+__device__
+inline
+int __any(int predicate) {
+ return __ockl_wfany_i32(predicate);
+}
+
+// XXX from llvm/include/llvm/IR/InstrTypes.h
+#define ICMP_NE 33
+
+__device__
+inline
+unsigned long long int __ballot(int predicate) {
+ return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE);
+}
+
+__device__
+inline
+unsigned long long int __ballot64(int predicate) {
+ return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE);
+}
+
+// hip.amdgcn.bc - lanemask
+__device__
+inline
+uint64_t __lanemask_gt()
+{
+ uint32_t lane = __ockl_lane_u32();
+ if (lane == 63)
+ return 0;
+ uint64_t ballot = __ballot64(1);
+ uint64_t mask = (~((uint64_t)0)) << (lane + 1);
+ return mask & ballot;
+}
+
+__device__
+inline
+uint64_t __lanemask_lt()
+{
+ uint32_t lane = __ockl_lane_u32();
+ int64_t ballot = __ballot64(1);
+ uint64_t mask = ((uint64_t)1 << lane) - (uint64_t)1;
+ return mask & ballot;
+}
+
+__device__
+inline
+uint64_t __lanemask_eq()
+{
+ uint32_t lane = __ockl_lane_u32();
+ int64_t mask = ((uint64_t)1 << lane);
+ return mask;
+}
+
+
+__device__ inline void* __local_to_generic(void* p) { return p; }
+
+#ifdef __HIP_DEVICE_COMPILE__
+__device__
+inline
+void* __get_dynamicgroupbaseptr()
+{
+ // Get group segment base pointer.
+ return (char*)__local_to_generic((void*)__to_local(__builtin_amdgcn_groupstaticsize()));
+}
+#else
+__device__
+void* __get_dynamicgroupbaseptr();
+#endif // __HIP_DEVICE_COMPILE__
+
+__device__
+inline
+void *__amdgcn_get_dynamicgroupbaseptr() {
+ return __get_dynamicgroupbaseptr();
+}
+
+// Memory Fence Functions
+__device__
+inline
+static void __threadfence()
+{
+ __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+}
+
+__device__
+inline
+static void __threadfence_block()
+{
+ __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+}
+
+__device__
+inline
+static void __threadfence_system()
+{
+ __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+}
+__device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) {
+ if (flags) {
+ __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
+ __builtin_amdgcn_s_barrier();
+ __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
+ } else {
+ __builtin_amdgcn_s_barrier();
+ }
+}
+
+__device__
+inline
+static void __barrier(int n)
+{
+ __work_group_barrier((__cl_mem_fence_flags)n);
+}
+
+__device__
+inline
+__attribute__((convergent))
+void __syncthreads()
+{
+ __barrier(__CLK_LOCAL_MEM_FENCE);
+}
+
+__device__
+inline
+__attribute__((convergent))
+int __syncthreads_count(int predicate)
+{
+ return __ockl_wgred_add_i32(!!predicate);
+}
+
+__device__
+inline
+__attribute__((convergent))
+int __syncthreads_and(int predicate)
+{
+ return __ockl_wgred_and_i32(!!predicate);
+}
+
+__device__
+inline
+__attribute__((convergent))
+int __syncthreads_or(int predicate)
+{
+ return __ockl_wgred_or_i32(!!predicate);
+}
+
+// hip.amdgcn.bc - device routine
+/*
+ HW_ID Register bit structure for RDNA2 & RDNA3
+ WAVE_ID 4:0 Wave id within the SIMD.
+ SIMD_ID 9:8 SIMD_ID within the WGP: [0] = row, [1] = column.
+ WGP_ID 13:10 Physical WGP ID.
+ SA_ID 16 Shader Array ID
+ SE_ID 20:18 Shader Engine the wave is assigned to for gfx11
+ SE_ID 19:18 Shader Engine the wave is assigned to for gfx10
+ DP_RATE 31:29 Number of double-precision float units per SIMD
+
+ HW_ID Register bit structure for GCN and CDNA
+ WAVE_ID 3:0 Wave buffer slot number. 0-9.
+ SIMD_ID 5:4 SIMD which the wave is assigned to within the CU.
+ PIPE_ID 7:6 Pipeline from which the wave was dispatched.
+ CU_ID 11:8 Compute Unit the wave is assigned to.
+ SH_ID 12 Shader Array (within an SE) the wave is assigned to.
+ SE_ID 15:13 Shader Engine the wave is assigned to for gfx908, gfx90a, gfx940-942
+ 14:13 Shader Engine the wave is assigned to for Vega.
+ TG_ID 19:16 Thread-group ID
+ VM_ID 23:20 Virtual Memory ID
+ QUEUE_ID 26:24 Queue from which this wave was dispatched.
+ STATE_ID 29:27 State ID (graphics only, not compute).
+ ME_ID 31:30 Micro-engine ID.
+
+ XCC_ID Register bit structure for gfx940
+ XCC_ID 3:0 XCC the wave is assigned to.
+ */
+
+#if (defined (__GFX10__) || defined (__GFX11__))
+ #define HW_ID 23
+#else
+ #define HW_ID 4
+#endif
+
+#if (defined(__GFX10__) || defined(__GFX11__))
+ #define HW_ID_WGP_ID_SIZE 4
+ #define HW_ID_WGP_ID_OFFSET 10
+#else
+ #define HW_ID_CU_ID_SIZE 4
+ #define HW_ID_CU_ID_OFFSET 8
+#endif
+
+#if (defined(__gfx908__) || defined(__gfx90a__) || \
+ defined(__GFX11__))
+ #define HW_ID_SE_ID_SIZE 3
+#else //4 SEs/XCC for gfx940-942
+ #define HW_ID_SE_ID_SIZE 2
+#endif
+#if (defined(__GFX10__) || defined(__GFX11__))
+ #define HW_ID_SE_ID_OFFSET 18
+ #define HW_ID_SA_ID_OFFSET 16
+ #define HW_ID_SA_ID_SIZE 1
+#else
+ #define HW_ID_SE_ID_OFFSET 13
+#endif
+
+#if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__))
+ #define XCC_ID 20
+ #define XCC_ID_XCC_ID_SIZE 4
+ #define XCC_ID_XCC_ID_OFFSET 0
+#endif
+
+#if (!defined(__HIP_NO_IMAGE_SUPPORT) && \
+ (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)))
+ #define __HIP_NO_IMAGE_SUPPORT 1
+#endif
+
+/*
+ Encoding of parameter bitmask
+ HW_ID 5:0 HW_ID
+ OFFSET 10:6 Range: 0..31
+ SIZE 15:11 Range: 1..32
+ */
+
+#define GETREG_IMMED(SZ,OFF,REG) (((SZ) << 11) | ((OFF) << 6) | (REG))
+
+/*
+ __smid returns the wave's assigned Compute Unit and Shader Engine.
+ The Compute Unit, CU_ID returned in bits 3:0, and Shader Engine, SE_ID in bits 5:4.
+ Note: the results vary over time.
+ SZ minus 1 since SIZE is 1-based.
+*/
+__device__
+inline
+unsigned __smid(void)
+{
+ unsigned se_id = __builtin_amdgcn_s_getreg(
+ GETREG_IMMED(HW_ID_SE_ID_SIZE-1, HW_ID_SE_ID_OFFSET, HW_ID));
+ #if (defined(__GFX10__) || defined(__GFX11__))
+ unsigned wgp_id = __builtin_amdgcn_s_getreg(
+ GETREG_IMMED(HW_ID_WGP_ID_SIZE - 1, HW_ID_WGP_ID_OFFSET, HW_ID));
+ unsigned sa_id = __builtin_amdgcn_s_getreg(
+ GETREG_IMMED(HW_ID_SA_ID_SIZE - 1, HW_ID_SA_ID_OFFSET, HW_ID));
+ #else
+ #if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__))
+ unsigned xcc_id = __builtin_amdgcn_s_getreg(
+ GETREG_IMMED(XCC_ID_XCC_ID_SIZE - 1, XCC_ID_XCC_ID_OFFSET, XCC_ID));
+ #endif
+ unsigned cu_id = __builtin_amdgcn_s_getreg(
+ GETREG_IMMED(HW_ID_CU_ID_SIZE - 1, HW_ID_CU_ID_OFFSET, HW_ID));
+ #endif
+ #if (defined(__GFX10__) || defined(__GFX11__))
+ unsigned temp = se_id;
+ temp = (temp << HW_ID_SA_ID_SIZE) | sa_id;
+ temp = (temp << HW_ID_WGP_ID_SIZE) | wgp_id;
+ return temp;
+ //TODO : CU Mode impl
+ #elif (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__))
+ unsigned temp = xcc_id;
+ temp = (temp << HW_ID_SE_ID_SIZE) | se_id;
+ temp = (temp << HW_ID_CU_ID_SIZE) | cu_id;
+ return temp;
+ #else
+ return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
+ #endif
+}
+
+/**
+ * Map HIP_DYNAMIC_SHARED to "extern __shared__" for compatibility with old HIP applications
+ * To be removed in a future release.
+ */
+#define HIP_DYNAMIC_SHARED(type, var) extern __shared__ type var[];
+#define HIP_DYNAMIC_SHARED_ATTRIBUTE
+
+#endif //defined(__clang__) && defined(__HIP__)
+
+
+// loop unrolling
+static inline __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size) {
+ auto dstPtr = static_cast(dst);
+ auto srcPtr = static_cast(src);
+
+ while (size >= 4u) {
+ dstPtr[0] = srcPtr[0];
+ dstPtr[1] = srcPtr[1];
+ dstPtr[2] = srcPtr[2];
+ dstPtr[3] = srcPtr[3];
+
+ size -= 4u;
+ srcPtr += 4u;
+ dstPtr += 4u;
+ }
+ switch (size) {
+ case 3:
+ dstPtr[2] = srcPtr[2];
+ case 2:
+ dstPtr[1] = srcPtr[1];
+ case 1:
+ dstPtr[0] = srcPtr[0];
+ }
+
+ return dst;
+}
+
+static inline __device__ void* __hip_hc_memset(void* dst, unsigned char val, size_t size) {
+ auto dstPtr = static_cast(dst);
+
+ while (size >= 4u) {
+ dstPtr[0] = val;
+ dstPtr[1] = val;
+ dstPtr[2] = val;
+ dstPtr[3] = val;
+
+ size -= 4u;
+ dstPtr += 4u;
+ }
+ switch (size) {
+ case 3:
+ dstPtr[2] = val;
+ case 2:
+ dstPtr[1] = val;
+ case 1:
+ dstPtr[0] = val;
+ }
+
+ return dst;
+}
+#ifndef __OPENMP_AMDGCN__
+static inline __device__ void* memcpy(void* dst, const void* src, size_t size) {
+ return __hip_hc_memcpy(dst, src, size);
+}
+
+static inline __device__ void* memset(void* ptr, int val, size_t size) {
+ unsigned char val8 = static_cast(val);
+ return __hip_hc_memset(ptr, val8, size);
+}
+#endif // !__OPENMP_AMDGCN__
+
+#endif
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_atomic.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_atomic.h
new file mode 100644
index 0000000000000000000000000000000000000000..ef719f3713c6546bc643842f33a16c65f9bd1bc6
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_atomic.h
@@ -0,0 +1,1612 @@
+/*
+Copyright (c) 2015 - Present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+
+#if !defined(__HIPCC_RTC__)
+#include "amd_device_functions.h"
+#endif
+
+#if __has_builtin(__hip_atomic_compare_exchange_strong)
+
+template struct Cond_t;
+
+template struct Cond_t { using type = T; };
+template struct Cond_t { using type = F; };
+
+#if !__HIP_DEVICE_COMPILE__
+//TODO: Remove this after compiler pre-defines the following Macros.
+#define __HIP_MEMORY_SCOPE_SINGLETHREAD 1
+#define __HIP_MEMORY_SCOPE_WAVEFRONT 2
+#define __HIP_MEMORY_SCOPE_WORKGROUP 3
+#define __HIP_MEMORY_SCOPE_AGENT 4
+#define __HIP_MEMORY_SCOPE_SYSTEM 5
+#endif
+
+#if !defined(__HIPCC_RTC__)
+#include "amd_hip_unsafe_atomics.h"
+#endif
+
+// Atomic expanders
+template<
+ int mem_order = __ATOMIC_SEQ_CST,
+ int mem_scope= __HIP_MEMORY_SCOPE_SYSTEM,
+ typename T,
+ typename Op,
+ typename F>
+inline
+__attribute__((always_inline, device))
+T hip_cas_expander(T* p, T x, Op op, F f) noexcept
+{
+ using FP = __attribute__((address_space(0))) const void*;
+
+ __device__
+ extern bool is_shared_workaround(FP) asm("llvm.amdgcn.is.shared");
+
+ if (is_shared_workaround((FP)p))
+ return f();
+
+ using U = typename Cond_t<
+ sizeof(T) == sizeof(unsigned int), unsigned int, unsigned long long>::type;
+
+ auto q = reinterpret_cast(p);
+
+ U tmp0{__hip_atomic_load(q, mem_order, mem_scope)};
+ U tmp1;
+ do {
+ tmp1 = tmp0;
+
+ op(reinterpret_cast(tmp1), x);
+ } while (!__hip_atomic_compare_exchange_strong(q, &tmp0, tmp1, mem_order,
+ mem_order, mem_scope));
+
+ return reinterpret_cast(tmp0);
+}
+
+template<
+ int mem_order = __ATOMIC_SEQ_CST,
+ int mem_scope= __HIP_MEMORY_SCOPE_SYSTEM,
+ typename T,
+ typename Cmp,
+ typename F>
+inline
+__attribute__((always_inline, device))
+T hip_cas_extrema_expander(T* p, T x, Cmp cmp, F f) noexcept
+{
+ using FP = __attribute__((address_space(0))) const void*;
+
+ __device__
+ extern bool is_shared_workaround(FP) asm("llvm.amdgcn.is.shared");
+
+ if (is_shared_workaround((FP)p))
+ return f();
+
+ using U = typename Cond_t<
+ sizeof(T) == sizeof(unsigned int), unsigned int, unsigned long long>::type;
+
+ auto q = reinterpret_cast(p);
+
+ U tmp{__hip_atomic_load(q, mem_order, mem_scope)};
+ while (cmp(x, reinterpret_cast(tmp)) &&
+ !__hip_atomic_compare_exchange_strong(q, &tmp, x, mem_order, mem_order,
+ mem_scope));
+
+ return reinterpret_cast(tmp);
+}
+
+__device__
+inline
+int atomicCAS(int* address, int compare, int val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ return compare;
+}
+
+__device__
+inline
+int atomicCAS_system(int* address, int compare, int val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ return compare;
+}
+
+__device__
+inline
+unsigned int atomicCAS(unsigned int* address, unsigned int compare, unsigned int val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ return compare;
+}
+
+__device__
+inline
+unsigned int atomicCAS_system(unsigned int* address, unsigned int compare, unsigned int val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ return compare;
+}
+
+__device__
+inline
+unsigned long atomicCAS(unsigned long* address, unsigned long compare, unsigned long val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ return compare;
+}
+
+__device__
+inline
+unsigned long atomicCAS_system(unsigned long* address, unsigned long compare, unsigned long val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ return compare;
+}
+
+__device__
+inline
+unsigned long long atomicCAS(unsigned long long* address, unsigned long long compare,
+ unsigned long long val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ return compare;
+}
+
+__device__
+inline
+unsigned long long atomicCAS_system(unsigned long long* address, unsigned long long compare,
+ unsigned long long val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ return compare;
+}
+
+__device__
+inline
+float atomicCAS(float* address, float compare, float val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ return compare;
+}
+
+__device__
+inline
+float atomicCAS_system(float* address, float compare, float val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ return compare;
+}
+
+__device__
+inline
+double atomicCAS(double* address, double compare, double val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ return compare;
+}
+
+__device__
+inline
+double atomicCAS_system(double* address, double compare, double val) {
+ __hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ return compare;
+}
+
+__device__
+inline
+int atomicAdd(int* address, int val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+int atomicAdd_system(int* address, int val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned int atomicAdd(unsigned int* address, unsigned int val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned int atomicAdd_system(unsigned int* address, unsigned int val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long atomicAdd(unsigned long* address, unsigned long val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicAdd_system(unsigned long* address, unsigned long val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long long atomicAdd(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long long atomicAdd_system(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+float atomicAdd(float* address, float val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicAdd(address, val);
+#else
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif
+}
+
+__device__
+inline
+float atomicAdd_system(float* address, float val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+#if !defined(__HIPCC_RTC__)
+DEPRECATED("use atomicAdd instead")
+#endif // !defined(__HIPCC_RTC__)
+__device__
+inline
+void atomicAddNoRet(float* address, float val)
+{
+ __ockl_atomic_add_noret_f32(address, val);
+}
+
+__device__
+inline
+double atomicAdd(double* address, double val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicAdd(address, val);
+#else
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif
+}
+
+__device__
+inline
+double atomicAdd_system(double* address, double val) {
+ return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+int atomicSub(int* address, int val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+int atomicSub_system(int* address, int val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned int atomicSub(unsigned int* address, unsigned int val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned int atomicSub_system(unsigned int* address, unsigned int val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long atomicSub(unsigned long* address, unsigned long val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicSub_system(unsigned long* address, unsigned long val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long long atomicSub(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long long atomicSub_system(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+float atomicSub(float* address, float val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicAdd(address, -val);
+#else
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif
+}
+
+__device__
+inline
+float atomicSub_system(float* address, float val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+double atomicSub(double* address, double val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicAdd(address, -val);
+#else
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif
+}
+
+__device__
+inline
+double atomicSub_system(double* address, double val) {
+ return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+int atomicExch(int* address, int val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+int atomicExch_system(int* address, int val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned int atomicExch(unsigned int* address, unsigned int val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned int atomicExch_system(unsigned int* address, unsigned int val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long atomicExch(unsigned long* address, unsigned long val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long atomicExch_system(unsigned long* address, unsigned long val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+unsigned long long atomicExch(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+unsigned long long atomicExch_system(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+float atomicExch(float* address, float val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+float atomicExch_system(float* address, float val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+double atomicExch(double* address, double val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+}
+
+__device__
+inline
+double atomicExch_system(double* address, double val) {
+ return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+__device__
+inline
+int atomicMin(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](int x, int y) { return x < y; }, [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicMin_system(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](int x, int y) { return x < y; }, [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicMin(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned int x, unsigned int y) { return x < y; }, [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+
+}
+
+__device__
+inline
+unsigned int atomicMin_system(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned int x, unsigned int y) { return x < y; }, [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicMin(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long x, unsigned long y) { return x < y; },
+ [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicMin_system(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address,
+ val,
+ [](unsigned long x, unsigned long y) { return x < y; },
+ [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicMin(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long long x, unsigned long long y) { return x < y; },
+ [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicMin_system(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address,
+ val,
+ [](unsigned long long x, unsigned long long y) { return x < y; },
+ [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+long long atomicMin(long long* address, long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](long long x, long long y) { return x < y; },
+ [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+long long atomicMin_system(long long* address, long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](long long x, long long y) { return x < y; },
+ [=]() {
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+float atomicMin(float* addr, float val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicMin(addr, val);
+#else
+ #if __has_builtin(__hip_atomic_load) && \
+ __has_builtin(__hip_atomic_compare_exchange_strong)
+ float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ bool done = false;
+ while (!done && value > val) {
+ done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ }
+ return value;
+ #else
+ unsigned int *uaddr = (unsigned int *)addr;
+ unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+ bool done = false;
+ while (!done && __uint_as_float(value) > val) {
+ done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+ }
+ return __uint_as_float(value);
+ #endif
+#endif
+}
+
+__device__
+inline
+float atomicMin_system(float* address, float val) {
+ unsigned int* uaddr { reinterpret_cast(address) };
+ #if __has_builtin(__hip_atomic_load)
+ unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+ #else
+ unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+ #endif
+ float value = __uint_as_float(tmp);
+
+ while (val < value) {
+ value = atomicCAS_system(address, value, val);
+ }
+
+ return value;
+}
+
+__device__
+inline
+double atomicMin(double* addr, double val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicMin(addr, val);
+#else
+ #if __has_builtin(__hip_atomic_load) && \
+ __has_builtin(__hip_atomic_compare_exchange_strong)
+ double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ bool done = false;
+ while (!done && value > val) {
+ done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ }
+ return value;
+ #else
+ unsigned long long *uaddr = (unsigned long long *)addr;
+ unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+ bool done = false;
+ while (!done && __longlong_as_double(value) > val) {
+ done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+ }
+ return __longlong_as_double(value);
+ #endif
+#endif
+}
+
+__device__
+inline
+double atomicMin_system(double* address, double val) {
+ unsigned long long* uaddr { reinterpret_cast(address) };
+ #if __has_builtin(__hip_atomic_load)
+ unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+ #else
+ unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+ #endif
+ double value = __longlong_as_double(tmp);
+
+ while (val < value) {
+ value = atomicCAS_system(address, value, val);
+ }
+
+ return value;
+}
+
+__device__
+inline
+int atomicMax(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](int x, int y) { return y < x; }, [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicMax_system(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](int x, int y) { return y < x; }, [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicMax(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned int x, unsigned int y) { return y < x; }, [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicMax_system(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned int x, unsigned int y) { return y < x; }, [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicMax(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long x, unsigned long y) { return y < x; },
+ [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicMax_system(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address,
+ val,
+ [](unsigned long x, unsigned long y) { return y < x; },
+ [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicMax(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long long x, unsigned long long y) { return y < x; },
+ [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicMax_system(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address,
+ val,
+ [](unsigned long long x, unsigned long long y) { return y < x; },
+ [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+long long atomicMax(long long* address, long long val) {
+ #if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](long long x, long long y) { return y < x; },
+ [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+long long atomicMax_system(long long* address, long long val) {
+#if defined(__gfx941__)
+ return hip_cas_extrema_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](long long x, long long y) { return y < x; },
+ [=]() {
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+float atomicMax(float* addr, float val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicMax(addr, val);
+#else
+ #if __has_builtin(__hip_atomic_load) && \
+ __has_builtin(__hip_atomic_compare_exchange_strong)
+ float value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ bool done = false;
+ while (!done && value < val) {
+ done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ }
+ return value;
+ #else
+ unsigned int *uaddr = (unsigned int *)addr;
+ unsigned int value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+ bool done = false;
+ while (!done && __uint_as_float(value) < val) {
+ done = __atomic_compare_exchange_n(uaddr, &value, __float_as_uint(val), false,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+ }
+ return __uint_as_float(value);
+ #endif
+#endif
+}
+
+__device__
+inline
+float atomicMax_system(float* address, float val) {
+ unsigned int* uaddr { reinterpret_cast(address) };
+ #if __has_builtin(__hip_atomic_load)
+ unsigned int tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+ #else
+ unsigned int tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+ #endif
+ float value = __uint_as_float(tmp);
+
+ while (value < val) {
+ value = atomicCAS_system(address, value, val);
+ }
+
+ return value;
+}
+
+__device__
+inline
+double atomicMax(double* addr, double val) {
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicMax(addr, val);
+#else
+ #if __has_builtin(__hip_atomic_load) && \
+ __has_builtin(__hip_atomic_compare_exchange_strong)
+ double value = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ bool done = false;
+ while (!done && value < val) {
+ done = __hip_atomic_compare_exchange_strong(addr, &value, val,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ }
+ return value;
+ #else
+ unsigned long long *uaddr = (unsigned long long *)addr;
+ unsigned long long value = __atomic_load_n(uaddr, __ATOMIC_RELAXED);
+ bool done = false;
+ while (!done && __longlong_as_double(value) < val) {
+ done = __atomic_compare_exchange_n(uaddr, &value, __double_as_longlong(val), false,
+ __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+ }
+ return __longlong_as_double(value);
+ #endif
+#endif
+}
+
+__device__
+inline
+double atomicMax_system(double* address, double val) {
+ unsigned long long* uaddr { reinterpret_cast(address) };
+ #if __has_builtin(__hip_atomic_load)
+ unsigned long long tmp {__hip_atomic_load(uaddr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM)};
+ #else
+ unsigned long long tmp {__atomic_load_n(uaddr, __ATOMIC_RELAXED)};
+ #endif
+ double value = __longlong_as_double(tmp);
+
+ while (value < val) {
+ value = atomicCAS_system(address, value, val);
+ }
+
+ return value;
+}
+
+__device__
+inline
+unsigned int atomicInc(unsigned int* address, unsigned int val)
+{
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned int& x, unsigned int y) { x = (x >= y) ? 0 : (x + 1); },
+ [=]() {
+ return
+ __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
+ });
+#else
+ return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
+#endif // __gfx941__
+
+}
+
+__device__
+inline
+unsigned int atomicDec(unsigned int* address, unsigned int val)
+{
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned int& x, unsigned int y) { x = (!x || x > y) ? y : (x - 1); },
+ [=]() {
+ return
+ __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
+ });
+#else
+ return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
+#endif // __gfx941__
+
+}
+
+__device__
+inline
+int atomicAnd(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](int& x, int y) { x &= y; }, [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicAnd_system(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](int& x, int y) { x &= y; }, [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicAnd(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned int& x, unsigned int y) { x &= y; }, [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicAnd_system(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned int& x, unsigned int y) { x &= y; }, [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicAnd(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned long& x, unsigned long y) { x &= y; }, [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicAnd_system(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned long& x, unsigned long y) { x &= y; }, [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicAnd(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long long& x, unsigned long long y) { x &= y; },
+ [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicAnd_system(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address,
+ val,
+ [](unsigned long long& x, unsigned long long y) { x &= y; },
+ [=]() {
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicOr(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](int& x, int y) { x |= y; }, [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicOr_system(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](int& x, int y) { x |= y; }, [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicOr(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned int& x, unsigned int y) { x |= y; }, [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicOr_system(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned int& x, unsigned int y) { x |= y; }, [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicOr(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned long& x, unsigned long y) { x |= y; }, [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicOr_system(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned long& x, unsigned long y) { x |= y; }, [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicOr(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long long& x, unsigned long long y) { x |= y; },
+ [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicOr_system(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address,
+ val,
+ [](unsigned long long& x, unsigned long long y) { x |= y; },
+ [=]() {
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicXor(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](int& x, int y) { x ^= y; }, [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+int atomicXor_system(int* address, int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](int& x, int y) { x ^= y; }, [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicXor(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned int& x, unsigned int y) { x ^= y; }, [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned int atomicXor_system(unsigned int* address, unsigned int val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned int& x, unsigned int y) { x ^= y; }, [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicXor(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address, val, [](unsigned long& x, unsigned long y) { x ^= y; }, [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long atomicXor_system(unsigned long* address, unsigned long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM>(
+ address, val, [](unsigned long& x, unsigned long y) { x ^= y; }, [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_SYSTEM);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicXor(unsigned long long* address, unsigned long long val) {
+#if defined(__gfx941__)
+ return hip_cas_expander<__ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT>(
+ address,
+ val,
+ [](unsigned long long& x, unsigned long long y) { x ^= y; },
+ [=]() {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED,
+ __HIP_MEMORY_SCOPE_AGENT);
+ });
+#else
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+#endif // __gfx941__
+}
+
+__device__
+inline
+unsigned long long atomicXor_system(unsigned long long* address, unsigned long long val) {
+ return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+}
+
+#else // __hip_atomic_compare_exchange_strong
+
+__device__
+inline
+int atomicCAS(int* address, int compare, int val)
+{
+ __atomic_compare_exchange_n(
+ address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+
+ return compare;
+}
+__device__
+inline
+unsigned int atomicCAS(
+ unsigned int* address, unsigned int compare, unsigned int val)
+{
+ __atomic_compare_exchange_n(
+ address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+
+ return compare;
+}
+__device__
+inline
+unsigned long long atomicCAS(
+ unsigned long long* address,
+ unsigned long long compare,
+ unsigned long long val)
+{
+ __atomic_compare_exchange_n(
+ address, &compare, val, false, __ATOMIC_RELAXED, __ATOMIC_RELAXED);
+
+ return compare;
+}
+
+__device__
+inline
+int atomicAdd(int* address, int val)
+{
+ return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicAdd(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicAdd(
+ unsigned long long* address, unsigned long long val)
+{
+ return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+float atomicAdd(float* address, float val)
+{
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicAdd(address, val);
+#else
+ return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+#endif
+}
+
+#if !defined(__HIPCC_RTC__)
+DEPRECATED("use atomicAdd instead")
+#endif // !defined(__HIPCC_RTC__)
+__device__
+inline
+void atomicAddNoRet(float* address, float val)
+{
+ __ockl_atomic_add_noret_f32(address, val);
+}
+
+__device__
+inline
+double atomicAdd(double* address, double val)
+{
+#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
+ return unsafeAtomicAdd(address, val);
+#else
+ return __atomic_fetch_add(address, val, __ATOMIC_RELAXED);
+#endif
+}
+
+__device__
+inline
+int atomicSub(int* address, int val)
+{
+ return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicSub(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_sub(address, val, __ATOMIC_RELAXED);
+}
+
+__device__
+inline
+int atomicExch(int* address, int val)
+{
+ return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicExch(unsigned int* address, unsigned int val)
+{
+ return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicExch(unsigned long long* address, unsigned long long val)
+{
+ return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+float atomicExch(float* address, float val)
+{
+ return __uint_as_float(__atomic_exchange_n(
+ reinterpret_cast(address),
+ __float_as_uint(val),
+ __ATOMIC_RELAXED));
+}
+
+__device__
+inline
+int atomicMin(int* address, int val)
+{
+ return __atomic_fetch_min(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicMin(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_min(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicMin(
+ unsigned long long* address, unsigned long long val)
+{
+ unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+ while (val < tmp) {
+ const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
+
+ if (tmp1 != tmp) { tmp = tmp1; continue; }
+
+ tmp = atomicCAS(address, tmp, val);
+ }
+
+ return tmp;
+}
+__device__ inline long long atomicMin(long long* address, long long val) {
+ long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+ while (val < tmp) {
+ const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
+
+ if (tmp1 != tmp) {
+ tmp = tmp1;
+ continue;
+ }
+
+ tmp = atomicCAS(address, tmp, val);
+ }
+ return tmp;
+}
+
+__device__
+inline
+int atomicMax(int* address, int val)
+{
+ return __atomic_fetch_max(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicMax(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_max(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicMax(
+ unsigned long long* address, unsigned long long val)
+{
+ unsigned long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+ while (tmp < val) {
+ const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
+
+ if (tmp1 != tmp) { tmp = tmp1; continue; }
+
+ tmp = atomicCAS(address, tmp, val);
+ }
+
+ return tmp;
+}
+__device__ inline long long atomicMax(long long* address, long long val) {
+ long long tmp{__atomic_load_n(address, __ATOMIC_RELAXED)};
+ while (tmp < val) {
+ const auto tmp1 = __atomic_load_n(address, __ATOMIC_RELAXED);
+
+ if (tmp1 != tmp) {
+ tmp = tmp1;
+ continue;
+ }
+
+ tmp = atomicCAS(address, tmp, val);
+ }
+ return tmp;
+}
+
+__device__
+inline
+unsigned int atomicInc(unsigned int* address, unsigned int val)
+{
+ return __builtin_amdgcn_atomic_inc32(address, val, __ATOMIC_RELAXED, "agent");
+}
+
+__device__
+inline
+unsigned int atomicDec(unsigned int* address, unsigned int val)
+{
+ return __builtin_amdgcn_atomic_dec32(address, val, __ATOMIC_RELAXED, "agent");
+}
+
+__device__
+inline
+int atomicAnd(int* address, int val)
+{
+ return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicAnd(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicAnd(
+ unsigned long long* address, unsigned long long val)
+{
+ return __atomic_fetch_and(address, val, __ATOMIC_RELAXED);
+}
+
+__device__
+inline
+int atomicOr(int* address, int val)
+{
+ return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicOr(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicOr(
+ unsigned long long* address, unsigned long long val)
+{
+ return __atomic_fetch_or(address, val, __ATOMIC_RELAXED);
+}
+
+__device__
+inline
+int atomicXor(int* address, int val)
+{
+ return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned int atomicXor(unsigned int* address, unsigned int val)
+{
+ return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
+}
+__device__
+inline
+unsigned long long atomicXor(
+ unsigned long long* address, unsigned long long val)
+{
+ return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
+}
+
+#endif // __hip_atomic_compare_exchange_strong
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_bf16.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_bf16.h
new file mode 100644
index 0000000000000000000000000000000000000000..204269a849c60e9b15ebec216d94ebd853783b55
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_bf16.h
@@ -0,0 +1,1337 @@
+/**
+ * MIT License
+ *
+ * Copyright (c) 2019 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+/**
+ * \file
+ * \brief hip_bf16.h provides struct for __hip_bfloat16 types
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT16 bfloat16 Precision Intrinsics
+ * This section describes hip_bfloat16 precision intrinsic functions.
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT16_ARITH Bfloat16 Arithmetic Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT16_COMP Bfloat16 Comparision Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT162_COMP Bfloat162 Comparision Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT162_ARITH Bfloat162 Arithmetic Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT16_CONV Bfloat16 Conversion Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT162_CONV Bfloat162 Conversion Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT16_MATH Bfloat16 Math Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+/**
+ * \defgroup HIP_INTRINSIC_BFLOAT162_MATH Bfloat162 Math Functions
+ * \ingroup HIP_INTRINSIC_BFLOAT16
+ * To use these functions, include the header file \p hip_bf16.h in your program.
+ */
+
+
+#ifndef _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_BF16_H_
+#define _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_BF16_H_
+
+#if !defined(__HIPCC_RTC__)
+#include
+#endif // !defined(__HIPCC_RTC__)
+
+#include "amd_hip_vector_types.h" // float2 etc
+#include "device_library_decls.h" // ocml conversion functions
+#include "math_fwd.h" // ocml device functions
+
+#if defined(__HIPCC_RTC__)
+#define __HOST_DEVICE__ __device__ static
+#else
+#include
+#include
+#include
+#define __HOST_DEVICE__ __host__ __device__ static inline
+#endif
+
+#define HIPRT_ONE_BF16 __float2bfloat16(1.0f)
+#define HIPRT_ZERO_BF16 __float2bfloat16(0.0f)
+#define HIPRT_INF_BF16 __ushort_as_bfloat16((unsigned short)0x7F80U)
+#define HIPRT_MAX_NORMAL_BF16 __ushort_as_bfloat16((unsigned short)0x7F7FU)
+#define HIPRT_MIN_DENORM_BF16 __ushort_as_bfloat16((unsigned short)0x0001U)
+#define HIPRT_NAN_BF16 __ushort_as_bfloat16((unsigned short)0x7FFFU)
+#define HIPRT_NEG_ZERO_BF16 __ushort_as_bfloat16((unsigned short)0x8000U)
+
+// Since we are using unsigned short to represent data in bfloat16, it can be of different sizes on
+// different machines. These naive checks should prevent some undefined behavior on systems which
+// have different sizes for basic types.
+#if !defined(__HIPCC_RTC__)
+static_assert(CHAR_BIT == 8, "byte size should be of 8 bits");
+#endif
+static_assert(sizeof(unsigned short) == 2, "size of unsigned short should be 2 bytes");
+
+/*! \brief Struct to represent a 16 bit brain floating point number. */
+struct __hip_bfloat16 {
+ unsigned short data;
+};
+
+/*! \brief Struct to represent two 16 bit brain floating point numbers. */
+struct __hip_bfloat162 {
+ __hip_bfloat16 x;
+ __hip_bfloat16 y;
+};
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_CONV
+ * \brief Converts bfloat16 to float
+ */
+__HOST_DEVICE__ inline float __bfloat162float(__hip_bfloat16 a) {
+ unsigned int uval = 0;
+ uval = a.data << 16;
+ union {
+ unsigned int u32;
+ float fp32;
+ } u = {uval};
+ return u.fp32;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_CONV
+ * \brief Converts float to bfloat16
+ */
+__HOST_DEVICE__ __hip_bfloat16 __float2bfloat16(float f) {
+ __hip_bfloat16 ret;
+ union {
+ float fp32;
+ unsigned int u32;
+ } u = {f};
+ if (~u.u32 & 0x7f800000) {
+ // When the exponent bits are not all 1s, then the value is zero, normal,
+ // or subnormal. We round the bfloat16 mantissa up by adding 0x7FFF, plus
+ // 1 if the least significant bit of the bfloat16 mantissa is 1 (odd).
+ // This causes the bfloat16's mantissa to be incremented by 1 if the 16
+ // least significant bits of the float mantissa are greater than 0x8000,
+ // or if they are equal to 0x8000 and the least significant bit of the
+ // bfloat16 mantissa is 1 (odd). This causes it to be rounded to even when
+ // the lower 16 bits are exactly 0x8000. If the bfloat16 mantissa already
+ // has the value 0x7f, then incrementing it causes it to become 0x00 and
+ // the exponent is incremented by one, which is the next higher FP value
+ // to the unrounded bfloat16 value. When the bfloat16 value is subnormal
+ // with an exponent of 0x00 and a mantissa of 0x7F, it may be rounded up
+ // to a normal value with an exponent of 0x01 and a mantissa of 0x00.
+ // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F,
+ // incrementing it causes it to become an exponent of 0xFF and a mantissa
+ // of 0x00, which is Inf, the next higher value to the unrounded value.
+ u.u32 += 0x7fff + ((u.u32 >> 16) & 1); // Round to nearest, round to even
+ } else if (u.u32 & 0xffff) {
+ // When all of the exponent bits are 1, the value is Inf or NaN.
+ // Inf is indicated by a zero mantissa. NaN is indicated by any nonzero
+ // mantissa bit. Quiet NaN is indicated by the most significant mantissa
+ // bit being 1. Signaling NaN is indicated by the most significant
+ // mantissa bit being 0 but some other bit(s) being 1. If any of the
+ // lower 16 bits of the mantissa are 1, we set the least significant bit
+ // of the bfloat16 mantissa, in order to preserve signaling NaN in case
+ // the bloat16's mantissa bits are all 0.
+ u.u32 |= 0x10000; // Preserve signaling NaN
+ }
+
+ ret.data = (u.u32 >> 16);
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Converts and moves bfloat162 to float2
+ */
+__HOST_DEVICE__ float2 __bfloat1622float2(const __hip_bfloat162 a) {
+ return float2{__bfloat162float(a.x), __bfloat162float(a.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Moves bfloat16 value to bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat162 __bfloat162bfloat162(const __hip_bfloat16 a) {
+ return __hip_bfloat162{a, a};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Reinterprets bits in a __hip_bfloat16 as a signed short integer
+ */
+__HOST_DEVICE__ short int __bfloat16_as_short(const __hip_bfloat16 h) { return (short)h.data; }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Reinterprets bits in a __hip_bfloat16 as an unsigned signed short integer
+ */
+__HOST_DEVICE__ unsigned short int __bfloat16_as_ushort(const __hip_bfloat16 h) { return h.data; }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Convert double to __hip_bfloat16
+ */
+__HOST_DEVICE__ __hip_bfloat16 __double2bfloat16(const double a) {
+ return __float2bfloat16((float)a);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Convert float2 to __hip_bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat162 __float22bfloat162_rn(const float2 a) {
+ return __hip_bfloat162{__float2bfloat16(a.x), __float2bfloat16(a.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Combine two __hip_bfloat16 to __hip_bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat162 __halves2bfloat162(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __hip_bfloat162{a, b};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Returns high 16 bits of __hip_bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat16 __high2bfloat16(const __hip_bfloat162 a) { return a.y; }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Returns high 16 bits of __hip_bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat162 __high2bfloat162(const __hip_bfloat162 a) {
+ return __hip_bfloat162{a.y, a.y};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Converts high 16 bits of __hip_bfloat162 to float and returns the result
+ */
+__HOST_DEVICE__ float __high2float(const __hip_bfloat162 a) { return __bfloat162float(a.y); }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Extracts high 16 bits from each and combines them
+ */
+__HOST_DEVICE__ __hip_bfloat162 __highs2bfloat162(const __hip_bfloat162 a,
+ const __hip_bfloat162 b) {
+ return __hip_bfloat162{a.y, b.y};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Returns low 16 bits of __hip_bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat16 __low2bfloat16(const __hip_bfloat162 a) { return a.x; }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Returns low 16 bits of __hip_bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat162 __low2bfloat162(const __hip_bfloat162 a) {
+ return __hip_bfloat162{a.x, a.x};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Converts low 16 bits of __hip_bfloat162 to float and returns the result
+ */
+__HOST_DEVICE__ float __low2float(const __hip_bfloat162 a) { return __bfloat162float(a.x); }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Swaps both halves
+ */
+__HOST_DEVICE__ __hip_bfloat162 __lowhigh2highlow(const __hip_bfloat162 a) {
+ return __hip_bfloat162{a.y, a.x};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Extracts low 16 bits from each and combines them
+ */
+__HOST_DEVICE__ __hip_bfloat162 __lows2bfloat162(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{a.x, b.x};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Reinterprets short int into a bfloat16
+ */
+__HOST_DEVICE__ __hip_bfloat16 __short_as_bfloat16(const short int a) {
+ return __hip_bfloat16{(unsigned short)a};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_CONV
+ * \brief Reinterprets unsigned short int into a bfloat16
+ */
+__HOST_DEVICE__ __hip_bfloat16 __ushort_as_bfloat16(const unsigned short int a) {
+ return __hip_bfloat16{a};
+}
+
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Adds two bfloat16 values
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hadd(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __float2bfloat16(__bfloat162float(a) + __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Subtracts two bfloat16 values
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hsub(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __float2bfloat16(__bfloat162float(a) - __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Divides two bfloat16 values
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hdiv(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __float2bfloat16(__bfloat162float(a) / __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Performs FMA of given bfloat16 values
+ */
+__device__ __hip_bfloat16 __hfma(const __hip_bfloat16 a, const __hip_bfloat16 b,
+ const __hip_bfloat16 c) {
+ return __float2bfloat16(
+ __ocml_fma_f32(__bfloat162float(a), __bfloat162float(b), __bfloat162float(c)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Multiplies two bfloat16 values
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hmul(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __float2bfloat16(__bfloat162float(a) * __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Negate a bfloat16 value
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hneg(const __hip_bfloat16 a) {
+ auto ret = a;
+ ret.data ^= 0x8000;
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Returns absolute of a bfloat16
+ */
+__HOST_DEVICE__ __hip_bfloat16 __habs(const __hip_bfloat16 a) {
+ auto ret = a;
+ ret.data &= 0x7FFF;
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Divides bfloat162 values
+ */
+__HOST_DEVICE__ __hip_bfloat162 __h2div(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{__float2bfloat16(__bfloat162float(a.x) / __bfloat162float(b.x)),
+ __float2bfloat16(__bfloat162float(a.y) / __bfloat162float(b.y))};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Returns absolute of a bfloat162
+ */
+__HOST_DEVICE__ __hip_bfloat162 __habs2(const __hip_bfloat162 a) {
+ return __hip_bfloat162{__habs(a.x), __habs(a.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Adds two bfloat162 values
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hadd2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{__hadd(a.x, b.x), __hadd(a.y, b.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Performs FMA of given bfloat162 values
+ */
+__device__ __hip_bfloat162 __hfma2(const __hip_bfloat162 a, const __hip_bfloat162 b,
+ const __hip_bfloat162 c) {
+ return __hip_bfloat162{__hfma(a.x, b.x, c.x), __hfma(a.y, b.y, c.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Multiplies two bfloat162 values
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hmul2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{__hmul(a.x, b.x), __hmul(a.y, b.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Converts a bfloat162 into negative
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hneg2(const __hip_bfloat162 a) {
+ return __hip_bfloat162{__hneg(a.x), __hneg(a.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Subtracts two bfloat162 values
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hsub2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{__hsub(a.x, b.x), __hsub(a.y, b.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to multiply two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator*(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hmul(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to multiply-assign two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16& operator*=(__hip_bfloat16& l, const __hip_bfloat16& r) {
+ l = __hmul(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to unary+ on a __hip_bfloat16 number
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator+(const __hip_bfloat16& l) { return l; }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to add two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator+(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hadd(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to negate a __hip_bfloat16 number
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator-(const __hip_bfloat16& l) { return __hneg(l); }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to subtract two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator-(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hsub(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to post increment a __hip_bfloat16 number
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator++(__hip_bfloat16& l, const int) {
+ auto ret = l;
+ l = __hadd(l, HIPRT_ONE_BF16);
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to pre increment a __hip_bfloat16 number
+ */
+__HOST_DEVICE__ __hip_bfloat16& operator++(__hip_bfloat16& l) {
+ l = __hadd(l, HIPRT_ONE_BF16);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to post decrement a __hip_bfloat16 number
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator--(__hip_bfloat16& l, const int) {
+ auto ret = l;
+ l = __hsub(l, HIPRT_ONE_BF16);
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to pre decrement a __hip_bfloat16 number
+ */
+__HOST_DEVICE__ __hip_bfloat16& operator--(__hip_bfloat16& l) {
+ l = __hsub(l, HIPRT_ONE_BF16);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to add-assign two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16& operator+=(__hip_bfloat16& l, const __hip_bfloat16& r) {
+ l = __hadd(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to subtract-assign two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16& operator-=(__hip_bfloat16& l, const __hip_bfloat16& r) {
+ l = __hsub(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to divide two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16 operator/(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hdiv(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_ARITH
+ * \brief Operator to divide-assign two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat16& operator/=(__hip_bfloat16& l, const __hip_bfloat16& r) {
+ l = __hdiv(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to multiply two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator*(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hmul2(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to multiply-assign two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162& operator*=(__hip_bfloat162& l, const __hip_bfloat162& r) {
+ l = __hmul2(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to unary+ on a __hip_bfloat162 number
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator+(const __hip_bfloat162& l) { return l; }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to add two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator+(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hadd2(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to negate a __hip_bfloat162 number
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator-(const __hip_bfloat162& l) { return __hneg2(l); }
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to subtract two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator-(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hsub2(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to post increment a __hip_bfloat162 number
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator++(__hip_bfloat162& l, const int) {
+ auto ret = l;
+ l = __hadd2(l, {HIPRT_ONE_BF16, HIPRT_ONE_BF16});
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to pre increment a __hip_bfloat162 number
+ */
+__HOST_DEVICE__ __hip_bfloat162& operator++(__hip_bfloat162& l) {
+ l = __hadd2(l, {HIPRT_ONE_BF16, HIPRT_ONE_BF16});
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to post decrement a __hip_bfloat162 number
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator--(__hip_bfloat162& l, const int) {
+ auto ret = l;
+ l = __hsub2(l, {HIPRT_ONE_BF16, HIPRT_ONE_BF16});
+ return ret;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to pre decrement a __hip_bfloat162 number
+ */
+__HOST_DEVICE__ __hip_bfloat162& operator--(__hip_bfloat162& l) {
+ l = __hsub2(l, {HIPRT_ONE_BF16, HIPRT_ONE_BF16});
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to add-assign two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162& operator+=(__hip_bfloat162& l, const __hip_bfloat162& r) {
+ l = __hadd2(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to subtract-assign two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162& operator-=(__hip_bfloat162& l, const __hip_bfloat162& r) {
+ l = __hsub2(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to divide two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162 operator/(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __h2div(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_ARITH
+ * \brief Operator to divide-assign two __hip_bfloat162 numbers
+ */
+__HOST_DEVICE__ __hip_bfloat162& operator/=(__hip_bfloat162& l, const __hip_bfloat162& r) {
+ l = __h2div(l, r);
+ return l;
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values
+ */
+__HOST_DEVICE__ bool __heq(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __bfloat162float(a) == __bfloat162float(b);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - unordered equal
+ */
+__HOST_DEVICE__ bool __hequ(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return !(__bfloat162float(a) < __bfloat162float(b)) &&
+ !(__bfloat162float(a) > __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - greater than
+ */
+__HOST_DEVICE__ bool __hgt(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __bfloat162float(a) > __bfloat162float(b);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - unordered greater than
+ */
+__HOST_DEVICE__ bool __hgtu(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return !(__bfloat162float(a) <= __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - greater than equal
+ */
+__HOST_DEVICE__ bool __hge(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __bfloat162float(a) >= __bfloat162float(b);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - unordered greater than equal
+ */
+__HOST_DEVICE__ bool __hgeu(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return !(__bfloat162float(a) < __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - not equal
+ */
+__HOST_DEVICE__ bool __hne(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __bfloat162float(a) != __bfloat162float(b);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - unordered not equal
+ */
+__HOST_DEVICE__ bool __hneu(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return !(__bfloat162float(a) == __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - return max
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hmax(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+#if __HIP_DEVICE_COMPILE__
+ return __float2bfloat16(__ocml_fmax_f32(__bfloat162float(a), __bfloat162float(b)));
+#else
+ return __float2bfloat16(std::max(__bfloat162float(a), __bfloat162float(b)));
+#endif
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - return min
+ */
+__HOST_DEVICE__ __hip_bfloat16 __hmin(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+#if __HIP_DEVICE_COMPILE__
+ return __float2bfloat16(__ocml_fmin_f32(__bfloat162float(a), __bfloat162float(b)));
+#else
+ return __float2bfloat16(std::min(__bfloat162float(a), __bfloat162float(b)));
+#endif
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - less than operator
+ */
+__HOST_DEVICE__ bool __hlt(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __bfloat162float(a) < __bfloat162float(b);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - unordered less than
+ */
+__HOST_DEVICE__ bool __hltu(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return !(__bfloat162float(a) >= __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - less than equal
+ */
+__HOST_DEVICE__ bool __hle(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return __bfloat162float(a) <= __bfloat162float(b);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Compare two bfloat162 values - unordered less than equal
+ */
+__HOST_DEVICE__ bool __hleu(const __hip_bfloat16 a, const __hip_bfloat16 b) {
+ return !(__bfloat162float(a) > __bfloat162float(b));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Checks if number is inf
+ */
+__HOST_DEVICE__ int __hisinf(const __hip_bfloat16 a) {
+ unsigned short sign = a.data & 0x8000U;
+#if __HIP_DEVICE_COMPILE__
+ int res = __ocml_isinf_f32(__bfloat162float(a));
+#else
+ int res = std::isinf(__bfloat162float(a)) ? 1 : 0;
+#endif
+ return (res == 0) ? res : ((sign != 0U) ? -res : res);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Checks if number is nan
+ */
+__HOST_DEVICE__ bool __hisnan(const __hip_bfloat16 a) {
+#if __HIP_DEVICE_COMPILE__
+ return __ocml_isnan_f32(__bfloat162float(a));
+#else
+ return std::isnan(__bfloat162float(a));
+#endif
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Checks if two numbers are equal
+ */
+__HOST_DEVICE__ bool __hbeq2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __heq(a.x, b.x) && __heq(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Checks if two numbers are equal - unordered
+ */
+__HOST_DEVICE__ bool __hbequ2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hequ(a.x, b.x) && __hequ(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a >= b
+ */
+__HOST_DEVICE__ bool __hbge2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hge(a.x, b.x) && __hge(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a >= b - unordered
+ */
+__HOST_DEVICE__ bool __hbgeu2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hgeu(a.x, b.x) && __hgeu(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a > b
+ */
+__HOST_DEVICE__ bool __hbgt2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hgt(a.x, b.x) && __hgt(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a > b - unordered
+ */
+__HOST_DEVICE__ bool __hbgtu2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hgtu(a.x, b.x) && __hgtu(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a <= b
+ */
+__HOST_DEVICE__ bool __hble2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hle(a.x, b.x) && __hle(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a <= b - unordered
+ */
+__HOST_DEVICE__ bool __hbleu2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hleu(a.x, b.x) && __hleu(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a < b
+ */
+__HOST_DEVICE__ bool __hblt2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hlt(a.x, b.x) && __hlt(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a < b - unordered
+ */
+__HOST_DEVICE__ bool __hbltu2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hltu(a.x, b.x) && __hltu(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a != b
+ */
+__HOST_DEVICE__ bool __hbne2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hne(a.x, b.x) && __hne(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a != b
+ */
+__HOST_DEVICE__ bool __hbneu2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hneu(a.x, b.x) && __hneu(a.y, b.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a != b, returns 1.0 if equal, otherwise 0.0
+ */
+__HOST_DEVICE__ __hip_bfloat162 __heq2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{{__heq(a.x, b.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__heq(a.y, b.y) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a >= b, returns 1.0 if greater than equal, otherwise 0.0
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hge2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{{__hge(a.x, b.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__hge(a.y, b.y) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a > b, returns 1.0 if greater than equal, otherwise 0.0
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hgt2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{{__hgt(a.x, b.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__hgt(a.y, b.y) ? HIPRT_ONE_BF16 : HIPRT_ONE_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a is NaN, returns 1.0 if NaN, otherwise 0.0
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hisnan2(const __hip_bfloat162 a) {
+ return __hip_bfloat162{{__hisnan(a.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__hisnan(a.y) ? HIPRT_ONE_BF16 : HIPRT_ONE_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a <= b, returns 1.0 if greater than equal, otherwise 0.0
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hle2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{{__hle(a.x, b.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__hle(a.y, b.y) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Check for a < b, returns 1.0 if greater than equal, otherwise 0.0
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hlt2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{{__hlt(a.x, b.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__hlt(a.y, b.y) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Returns max of two elements
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hmax2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{__hmax(a.x, b.x), __hmax(a.y, b.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Returns min of two elements
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hmin2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{__hmin(a.x, b.x), __hmin(a.y, b.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Checks for not equal to
+ */
+__HOST_DEVICE__ __hip_bfloat162 __hne2(const __hip_bfloat162 a, const __hip_bfloat162 b) {
+ return __hip_bfloat162{{__hne(a.x, b.x) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16},
+ {__hne(a.y, b.y) ? HIPRT_ONE_BF16 : HIPRT_ZERO_BF16}};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform an equal compare on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator==(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __heq(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform a not equal on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator!=(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hne(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform a less than on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator<(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hlt(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform a less than equal on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator<=(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hle(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform a greater than on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator>(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hgt(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform a greater than equal on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator>=(const __hip_bfloat16& l, const __hip_bfloat16& r) {
+ return __hge(l, r);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Operator to perform an equal compare on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator==(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __heq(l.x, r.x) && __heq(l.y, r.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Operator to perform a not equal on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator!=(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hne(l.x, r.x) || __hne(l.y, r.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Operator to perform a less than on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator<(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hlt(l.x, r.x) && __hlt(l.y, r.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Operator to perform a less than equal on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator<=(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hle(l.x, r.x) && __hle(l.y, r.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_COMP
+ * \brief Operator to perform a greater than on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator>(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hgt(l.x, r.x) && __hgt(l.y, r.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_COMP
+ * \brief Operator to perform a greater than equal on two __hip_bfloat16 numbers
+ */
+__HOST_DEVICE__ bool operator>=(const __hip_bfloat162& l, const __hip_bfloat162& r) {
+ return __hge(l.x, r.x) && __hge(l.y, r.y);
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate ceil of bfloat16
+ */
+__device__ __hip_bfloat16 hceil(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_ceil_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate cosine of bfloat16
+ */
+__device__ __hip_bfloat16 hcos(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_cos_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate exponential of bfloat16
+ */
+__device__ __hip_bfloat16 hexp(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_exp_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate exponential 10 of bfloat16
+ */
+__device__ __hip_bfloat16 hexp10(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_exp10_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate exponential 2 of bfloat16
+ */
+__device__ __hip_bfloat16 hexp2(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_exp2_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate floor of bfloat16
+ */
+__device__ __hip_bfloat16 hfloor(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_floor_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate natural log of bfloat16
+ */
+__device__ __hip_bfloat16 hlog(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_log_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate log 10 of bfloat16
+ */
+__device__ __hip_bfloat16 hlog10(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_log10_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate log 2 of bfloat16
+ */
+__device__ __hip_bfloat16 hlog2(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_log2_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate reciprocal
+ */
+__device__ __hip_bfloat16 hrcp(const __hip_bfloat16 h) {
+ return __float2bfloat16(1.0f / (__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Round to nearest int
+ */
+__device__ __hip_bfloat16 hrint(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_rint_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Reciprocal square root
+ */
+__device__ __hip_bfloat16 hrsqrt(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_rsqrt_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate sin of bfloat16
+ */
+__device__ __hip_bfloat16 hsin(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_sin_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate sqrt of bfloat16
+ */
+__device__ __hip_bfloat16 hsqrt(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_sqrt_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT16_MATH
+ * \brief Calculate truncate of bfloat16
+ */
+__device__ __hip_bfloat16 htrunc(const __hip_bfloat16 h) {
+ return __float2bfloat16(__ocml_trunc_f32(__bfloat162float(h)));
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate ceil of bfloat162
+ */
+__device__ __hip_bfloat162 h2ceil(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hceil(h.x), hceil(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate cosine of bfloat162
+ */
+__device__ __hip_bfloat162 h2cos(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hcos(h.x), hcos(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate exponential of bfloat162
+ */
+__device__ __hip_bfloat162 h2exp(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hexp(h.x), hexp(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate exponential 10 of bfloat162
+ */
+__device__ __hip_bfloat162 h2exp10(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hexp10(h.x), hexp10(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate exponential 2 of bfloat162
+ */
+__device__ __hip_bfloat162 h2exp2(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hexp2(h.x), hexp2(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate floor of bfloat162
+ */
+__device__ __hip_bfloat162 h2floor(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hfloor(h.x), hfloor(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate natural log of bfloat162
+ */
+__device__ __hip_bfloat162 h2log(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hlog(h.x), hlog(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate log 10 of bfloat162
+ */
+__device__ __hip_bfloat162 h2log10(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hlog10(h.x), hlog10(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate log 2 of bfloat162
+ */
+__device__ __hip_bfloat162 h2log2(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hlog2(h.x), hlog2(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate vector reciprocal
+ */
+__device__ __hip_bfloat162 h2rcp(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hrcp(h.x), hrcp(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate vector round to nearest int
+ */
+__device__ __hip_bfloat162 h2rint(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hrint(h.x), hrint(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate vector reciprocal square root
+ */
+__device__ __hip_bfloat162 h2rsqrt(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hrsqrt(h.x), hrsqrt(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate sin of bfloat162
+ */
+__device__ __hip_bfloat162 h2sin(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hsin(h.x), hsin(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate sqrt of bfloat162
+ */
+__device__ __hip_bfloat162 h2sqrt(const __hip_bfloat162 h) {
+ return __hip_bfloat162{hsqrt(h.x), hsqrt(h.y)};
+}
+
+/**
+ * \ingroup HIP_INTRINSIC_BFLOAT162_MATH
+ * \brief Calculate truncate of bfloat162
+ */
+__device__ __hip_bfloat162 h2trunc(const __hip_bfloat162 h) {
+ return __hip_bfloat162{htrunc(h.x), htrunc(h.y)};
+}
+#endif
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_bfloat16.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_bfloat16.h
new file mode 100644
index 0000000000000000000000000000000000000000..deb3bfb7e24baad162e307044538b2ffa88e2c60
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_bfloat16.h
@@ -0,0 +1,293 @@
+/**
+ * MIT License
+ *
+ * Copyright (c) 2019 - 2022 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+/*!\file
+ * \brief hip_bfloat16.h provides struct for hip_bfloat16 typedef
+ */
+
+#ifndef _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_BFLOAT16_H_
+#define _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_BFLOAT16_H_
+
+#include "host_defines.h"
+#if defined(__HIPCC_RTC__)
+ #define __HOST_DEVICE__ __device__
+#else
+ #define __HOST_DEVICE__ __host__ __device__
+#endif
+
+#if __cplusplus < 201103L || !defined(__HIPCC__)
+
+// If this is a C compiler, C++ compiler below C++11, or a host-only compiler, we only
+// include a minimal definition of hip_bfloat16
+
+#include
+/*! \brief Struct to represent a 16 bit brain floating point number. */
+typedef struct
+{
+ uint16_t data;
+} hip_bfloat16;
+
+#else // __cplusplus < 201103L || !defined(__HIPCC__)
+
+#include
+
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wshadow"
+struct hip_bfloat16
+{
+ __hip_uint16_t data;
+
+ enum truncate_t
+ {
+ truncate
+ };
+
+ __HOST_DEVICE__ hip_bfloat16() = default;
+
+ // round upper 16 bits of IEEE float to convert to bfloat16
+ explicit __HOST_DEVICE__ hip_bfloat16(float f)
+ : data(float_to_bfloat16(f))
+ {
+ }
+
+ explicit __HOST_DEVICE__ hip_bfloat16(float f, truncate_t)
+ : data(truncate_float_to_bfloat16(f))
+ {
+ }
+
+ // zero extend lower 16 bits of bfloat16 to convert to IEEE float
+ __HOST_DEVICE__ operator float() const
+ {
+ union
+ {
+ uint32_t int32;
+ float fp32;
+ } u = {uint32_t(data) << 16};
+ return u.fp32;
+ }
+
+ __HOST_DEVICE__ hip_bfloat16 &operator=(const float& f)
+ {
+ data = float_to_bfloat16(f);
+ return *this;
+ }
+
+ static __HOST_DEVICE__ hip_bfloat16 round_to_bfloat16(float f)
+ {
+ hip_bfloat16 output;
+ output.data = float_to_bfloat16(f);
+ return output;
+ }
+
+ static __HOST_DEVICE__ hip_bfloat16 round_to_bfloat16(float f, truncate_t)
+ {
+ hip_bfloat16 output;
+ output.data = truncate_float_to_bfloat16(f);
+ return output;
+ }
+
+private:
+ static __HOST_DEVICE__ __hip_uint16_t float_to_bfloat16(float f)
+ {
+ union
+ {
+ float fp32;
+ uint32_t int32;
+ } u = {f};
+ if(~u.int32 & 0x7f800000)
+ {
+ // When the exponent bits are not all 1s, then the value is zero, normal,
+ // or subnormal. We round the bfloat16 mantissa up by adding 0x7FFF, plus
+ // 1 if the least significant bit of the bfloat16 mantissa is 1 (odd).
+ // This causes the bfloat16's mantissa to be incremented by 1 if the 16
+ // least significant bits of the float mantissa are greater than 0x8000,
+ // or if they are equal to 0x8000 and the least significant bit of the
+ // bfloat16 mantissa is 1 (odd). This causes it to be rounded to even when
+ // the lower 16 bits are exactly 0x8000. If the bfloat16 mantissa already
+ // has the value 0x7f, then incrementing it causes it to become 0x00 and
+ // the exponent is incremented by one, which is the next higher FP value
+ // to the unrounded bfloat16 value. When the bfloat16 value is subnormal
+ // with an exponent of 0x00 and a mantissa of 0x7F, it may be rounded up
+ // to a normal value with an exponent of 0x01 and a mantissa of 0x00.
+ // When the bfloat16 value has an exponent of 0xFE and a mantissa of 0x7F,
+ // incrementing it causes it to become an exponent of 0xFF and a mantissa
+ // of 0x00, which is Inf, the next higher value to the unrounded value.
+ u.int32 += 0x7fff + ((u.int32 >> 16) & 1); // Round to nearest, round to even
+ }
+ else if(u.int32 & 0xffff)
+ {
+ // When all of the exponent bits are 1, the value is Inf or NaN.
+ // Inf is indicated by a zero mantissa. NaN is indicated by any nonzero
+ // mantissa bit. Quiet NaN is indicated by the most significant mantissa
+ // bit being 1. Signaling NaN is indicated by the most significant
+ // mantissa bit being 0 but some other bit(s) being 1. If any of the
+ // lower 16 bits of the mantissa are 1, we set the least significant bit
+ // of the bfloat16 mantissa, in order to preserve signaling NaN in case
+ // the bloat16's mantissa bits are all 0.
+ u.int32 |= 0x10000; // Preserve signaling NaN
+ }
+ return __hip_uint16_t(u.int32 >> 16);
+ }
+
+ // Truncate instead of rounding, preserving SNaN
+ static __HOST_DEVICE__ __hip_uint16_t truncate_float_to_bfloat16(float f)
+ {
+ union
+ {
+ float fp32;
+ uint32_t int32;
+ } u = {f};
+ return __hip_uint16_t(u.int32 >> 16) | (!(~u.int32 & 0x7f800000) && (u.int32 & 0xffff));
+ }
+};
+#pragma clang diagnostic pop
+
+typedef struct
+{
+ __hip_uint16_t data;
+} hip_bfloat16_public;
+
+static_assert(__hip_internal::is_standard_layout{},
+ "hip_bfloat16 is not a standard layout type, and thus is "
+ "incompatible with C.");
+
+static_assert(__hip_internal::is_trivial{},
+ "hip_bfloat16 is not a trivial type, and thus is "
+ "incompatible with C.");
+#if !defined(__HIPCC_RTC__)
+static_assert(sizeof(hip_bfloat16) == sizeof(hip_bfloat16_public)
+ && offsetof(hip_bfloat16, data) == offsetof(hip_bfloat16_public, data),
+ "internal hip_bfloat16 does not match public hip_bfloat16");
+
+inline std::ostream& operator<<(std::ostream& os, const hip_bfloat16& bf16)
+{
+ return os << float(bf16);
+}
+#endif
+
+inline __HOST_DEVICE__ hip_bfloat16 operator+(hip_bfloat16 a)
+{
+ return a;
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator-(hip_bfloat16 a)
+{
+ a.data ^= 0x8000;
+ return a;
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator+(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return hip_bfloat16(float(a) + float(b));
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator-(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return hip_bfloat16(float(a) - float(b));
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator*(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return hip_bfloat16(float(a) * float(b));
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator/(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return hip_bfloat16(float(a) / float(b));
+}
+inline __HOST_DEVICE__ bool operator<(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return float(a) < float(b);
+}
+inline __HOST_DEVICE__ bool operator==(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return float(a) == float(b);
+}
+inline __HOST_DEVICE__ bool operator>(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return b < a;
+}
+inline __HOST_DEVICE__ bool operator<=(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return !(a > b);
+}
+inline __HOST_DEVICE__ bool operator!=(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return !(a == b);
+}
+inline __HOST_DEVICE__ bool operator>=(hip_bfloat16 a, hip_bfloat16 b)
+{
+ return !(a < b);
+}
+inline __HOST_DEVICE__ hip_bfloat16& operator+=(hip_bfloat16& a, hip_bfloat16 b)
+{
+ return a = a + b;
+}
+inline __HOST_DEVICE__ hip_bfloat16& operator-=(hip_bfloat16& a, hip_bfloat16 b)
+{
+ return a = a - b;
+}
+inline __HOST_DEVICE__ hip_bfloat16& operator*=(hip_bfloat16& a, hip_bfloat16 b)
+{
+ return a = a * b;
+}
+inline __HOST_DEVICE__ hip_bfloat16& operator/=(hip_bfloat16& a, hip_bfloat16 b)
+{
+ return a = a / b;
+}
+inline __HOST_DEVICE__ hip_bfloat16& operator++(hip_bfloat16& a)
+{
+ return a += hip_bfloat16(1.0f);
+}
+inline __HOST_DEVICE__ hip_bfloat16& operator--(hip_bfloat16& a)
+{
+ return a -= hip_bfloat16(1.0f);
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator++(hip_bfloat16& a, int)
+{
+ hip_bfloat16 orig = a;
+ ++a;
+ return orig;
+}
+inline __HOST_DEVICE__ hip_bfloat16 operator--(hip_bfloat16& a, int)
+{
+ hip_bfloat16 orig = a;
+ --a;
+ return orig;
+}
+
+namespace std
+{
+ constexpr __HOST_DEVICE__ bool isinf(hip_bfloat16 a)
+ {
+ return !(~a.data & 0x7f80) && !(a.data & 0x7f);
+ }
+ constexpr __HOST_DEVICE__ bool isnan(hip_bfloat16 a)
+ {
+ return !(~a.data & 0x7f80) && +(a.data & 0x7f);
+ }
+ constexpr __HOST_DEVICE__ bool iszero(hip_bfloat16 a)
+ {
+ return !(a.data & 0x7fff);
+ }
+}
+
+#endif // __cplusplus < 201103L || !defined(__HIPCC__)
+
+#endif // _HIP_BFLOAT16_H_
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_common.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_common.h
new file mode 100644
index 0000000000000000000000000000000000000000..0c7dc51b503e5101abc365a2a62be471d00962e1
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_common.h
@@ -0,0 +1,32 @@
+/*
+Copyright (c) 2019 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy of
+this software and associated documentation files (the "Software"), to deal in
+the Software without restriction, including without limitation the rights to
+use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies
+of the Software, and to permit persons to whom the Software is furnished to do
+so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMMON_H
+#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMMON_H
+
+#if defined(__clang__) && defined(__HIP__)
+#define __HIP_CLANG_ONLY__ 1
+#else
+#define __HIP_CLANG_ONLY__ 0
+#endif
+
+#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMMON_H
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_complex.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_complex.h
new file mode 100644
index 0000000000000000000000000000000000000000..9ca29eafd206183a7de681aab31d033dea07f9f6
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_complex.h
@@ -0,0 +1,174 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* The header defines complex numbers and related functions*/
+
+#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
+#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
+
+#if !defined(__HIPCC_RTC__)
+#include "hip/amd_detail/amd_hip_vector_types.h"
+#endif
+
+#if defined(__HIPCC_RTC__)
+#define __HOST_DEVICE__ __device__
+#else
+#define __HOST_DEVICE__ __host__ __device__
+// TODO: Clang has a bug which allows device functions to call std functions
+// when std functions are introduced into default namespace by using statement.
+// math.h may be included after this bug is fixed.
+#if __cplusplus
+#include
+#else
+#include "math.h"
+#endif
+#endif // !defined(__HIPCC_RTC__)
+
+typedef float2 hipFloatComplex;
+
+__HOST_DEVICE__ static inline float hipCrealf(hipFloatComplex z) { return z.x; }
+
+__HOST_DEVICE__ static inline float hipCimagf(hipFloatComplex z) { return z.y; }
+
+__HOST_DEVICE__ static inline hipFloatComplex make_hipFloatComplex(float a, float b) {
+ hipFloatComplex z;
+ z.x = a;
+ z.y = b;
+ return z;
+}
+
+__HOST_DEVICE__ static inline hipFloatComplex hipConjf(hipFloatComplex z) {
+ hipFloatComplex ret;
+ ret.x = z.x;
+ ret.y = -z.y;
+ return ret;
+}
+
+__HOST_DEVICE__ static inline float hipCsqabsf(hipFloatComplex z) {
+ return z.x * z.x + z.y * z.y;
+}
+
+__HOST_DEVICE__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q) {
+ return make_hipFloatComplex(p.x + q.x, p.y + q.y);
+}
+
+__HOST_DEVICE__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q) {
+ return make_hipFloatComplex(p.x - q.x, p.y - q.y);
+}
+
+__HOST_DEVICE__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q) {
+ return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
+}
+
+__HOST_DEVICE__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q) {
+ float sqabs = hipCsqabsf(q);
+ hipFloatComplex ret;
+ ret.x = (p.x * q.x + p.y * q.y) / sqabs;
+ ret.y = (p.y * q.x - p.x * q.y) / sqabs;
+ return ret;
+}
+
+__HOST_DEVICE__ static inline float hipCabsf(hipFloatComplex z) { return sqrtf(hipCsqabsf(z)); }
+
+
+typedef double2 hipDoubleComplex;
+
+__HOST_DEVICE__ static inline double hipCreal(hipDoubleComplex z) { return z.x; }
+
+__HOST_DEVICE__ static inline double hipCimag(hipDoubleComplex z) { return z.y; }
+
+__HOST_DEVICE__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b) {
+ hipDoubleComplex z;
+ z.x = a;
+ z.y = b;
+ return z;
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipConj(hipDoubleComplex z) {
+ hipDoubleComplex ret;
+ ret.x = z.x;
+ ret.y = -z.y;
+ return ret;
+}
+
+__HOST_DEVICE__ static inline double hipCsqabs(hipDoubleComplex z) {
+ return z.x * z.x + z.y * z.y;
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q) {
+ return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q) {
+ return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q) {
+ return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q) {
+ double sqabs = hipCsqabs(q);
+ hipDoubleComplex ret;
+ ret.x = (p.x * q.x + p.y * q.y) / sqabs;
+ ret.y = (p.y * q.x - p.x * q.y) / sqabs;
+ return ret;
+}
+
+__HOST_DEVICE__ static inline double hipCabs(hipDoubleComplex z) { return sqrt(hipCsqabs(z)); }
+
+typedef hipFloatComplex hipComplex;
+
+__HOST_DEVICE__ static inline hipComplex make_hipComplex(float x, float y) {
+ return make_hipFloatComplex(x, y);
+}
+
+__HOST_DEVICE__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z) {
+ return make_hipFloatComplex((float)z.x, (float)z.y);
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z) {
+ return make_hipDoubleComplex((double)z.x, (double)z.y);
+}
+
+__HOST_DEVICE__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r) {
+ float real = (p.x * q.x) + r.x;
+ float imag = (q.x * p.y) + r.y;
+
+ real = -(p.y * q.y) + real;
+ imag = (p.x * q.y) + imag;
+
+ return make_hipComplex(real, imag);
+}
+
+__HOST_DEVICE__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q,
+ hipDoubleComplex r) {
+ double real = (p.x * q.x) + r.x;
+ double imag = (q.x * p.y) + r.y;
+
+ real = -(p.y * q.y) + real;
+ imag = (p.x * q.y) + imag;
+
+ return make_hipDoubleComplex(real, imag);
+}
+
+#endif //HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COMPLEX_H
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_cooperative_groups.h
new file mode 100644
index 0000000000000000000000000000000000000000..8b1a0c067db13ad77dc0c938a4327627a41a7ac3
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_cooperative_groups.h
@@ -0,0 +1,829 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/**
+ * @file amd_detail/hip_cooperative_groups.h
+ *
+ * @brief Device side implementation of `Cooperative Group` feature.
+ *
+ * Defines new types and device API wrappers related to `Cooperative Group`
+ * feature, which the programmer can directly use in his kernel(s) in order to
+ * make use of this feature.
+ */
+#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H
+#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H
+
+#if __cplusplus
+#if !defined(__HIPCC_RTC__)
+#include
+#endif
+
+namespace cooperative_groups {
+
+/** @brief The base type of all cooperative group types
+ *
+ * \details Holds the key properties of a constructed cooperative group types
+ * object, like the group type, its size, etc
+ *
+ * @note Cooperative groups feature is implemented on Linux, under developement
+ * on Windows.
+ */
+class thread_group {
+ protected:
+ uint32_t _type; // thread_group type
+ uint32_t _size; // total number of threads in the tread_group
+ uint64_t _mask; // Lanemask for coalesced and tiled partitioned group types,
+ // LSB represents lane 0, and MSB represents lane 63
+
+ // Construct a thread group, and set thread group type and other essential
+ // thread group properties. This generic thread group is directly constructed
+ // only when the group is supposed to contain only the calling the thread
+ // (throurh the API - `this_thread()`), and in all other cases, this thread
+ // group object is a sub-object of some other derived thread group object
+ __CG_QUALIFIER__ thread_group(internal::group_type type, uint32_t size = static_cast(0),
+ uint64_t mask = static_cast(0)) {
+ _type = type;
+ _size = size;
+ _mask = mask;
+ }
+
+ struct _tiled_info {
+ bool is_tiled;
+ unsigned int size;
+ unsigned int meta_group_rank;
+ unsigned int meta_group_size;
+ };
+
+ struct _coalesced_info {
+ lane_mask member_mask;
+ unsigned int size;
+ struct _tiled_info tiled_info;
+ } coalesced_info;
+
+ friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
+ unsigned int tile_size);
+ friend class thread_block;
+
+ public:
+ // Total number of threads in the thread group, and this serves the purpose
+ // for all derived cooperative group types since their `size` is directly
+ // saved during the construction
+ __CG_QUALIFIER__ uint32_t size() const { return _size; }
+ __CG_QUALIFIER__ unsigned int cg_type() const { return _type; }
+ // Rank of the calling thread within [0, size())
+ __CG_QUALIFIER__ uint32_t thread_rank() const;
+ // Is this cooperative group type valid?
+ __CG_QUALIFIER__ bool is_valid() const;
+ // synchronize the threads in the thread group
+ __CG_QUALIFIER__ void sync() const;
+};
+/**
+ *-------------------------------------------------------------------------------------------------
+ *-------------------------------------------------------------------------------------------------
+ * @defgroup CooperativeG Cooperative Groups
+ * @ingroup API
+ * @{
+ * This section describes the cooperative groups functions of HIP runtime API.
+ *
+ * The cooperative groups provides flexible thread parallel programming algorithms, threads
+ * cooperate and share data to perform collective computations.
+ *
+ * @note Cooperative groups feature is implemented on Linux, under developement
+ * on Windows.
+ *
+ */
+/** \brief The multi-grid cooperative group type
+ *
+ * \details Represents an inter-device cooperative group type where the
+ * participating threads within the group spans across multple
+ * devices, running the (same) kernel on these devices
+ * @note The multi-grid cooperative group type is implemented on Linux, under developement
+ * on Windows.
+ */
+class multi_grid_group : public thread_group {
+ // Only these friend functions are allowed to construct an object of this class
+ // and access its resources
+ friend __CG_QUALIFIER__ multi_grid_group this_multi_grid();
+
+ protected:
+ // Construct mutli-grid thread group (through the API this_multi_grid())
+ explicit __CG_QUALIFIER__ multi_grid_group(uint32_t size)
+ : thread_group(internal::cg_multi_grid, size) {}
+
+ public:
+ // Number of invocations participating in this multi-grid group. In other
+ // words, the number of GPUs
+ __CG_QUALIFIER__ uint32_t num_grids() { return internal::multi_grid::num_grids(); }
+ // Rank of this invocation. In other words, an ID number within the range
+ // [0, num_grids()) of the GPU, this kernel is running on
+ __CG_QUALIFIER__ uint32_t grid_rank() { return internal::multi_grid::grid_rank(); }
+ __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::multi_grid::thread_rank(); }
+ __CG_QUALIFIER__ bool is_valid() const { return internal::multi_grid::is_valid(); }
+ __CG_QUALIFIER__ void sync() const { internal::multi_grid::sync(); }
+};
+
+/** @brief User exposed API interface to construct multi-grid cooperative
+ * group type object - `multi_grid_group`
+ *
+ * \details User is not allowed to directly construct an object of type
+ * `multi_grid_group`. Instead, he should construct it through this
+ * API function
+ * @note This multi-grid cooperative API type is implemented on Linux, under developement
+ * on Windows.
+ */
+__CG_QUALIFIER__ multi_grid_group this_multi_grid() {
+ return multi_grid_group(internal::multi_grid::size());
+}
+
+/** @brief The grid cooperative group type
+ *
+ * \details Represents an inter-workgroup cooperative group type where the
+ * participating threads within the group spans across multiple
+ * workgroups running the (same) kernel on the same device
+ * @note This is implemented on Linux, under developement
+ * on Windows.
+ */
+class grid_group : public thread_group {
+ // Only these friend functions are allowed to construct an object of this class
+ // and access its resources
+ friend __CG_QUALIFIER__ grid_group this_grid();
+
+ protected:
+ // Construct grid thread group (through the API this_grid())
+ explicit __CG_QUALIFIER__ grid_group(uint32_t size) : thread_group(internal::cg_grid, size) {}
+
+ public:
+ __CG_QUALIFIER__ uint32_t thread_rank() const { return internal::grid::thread_rank(); }
+ __CG_QUALIFIER__ bool is_valid() const { return internal::grid::is_valid(); }
+ __CG_QUALIFIER__ void sync() const { internal::grid::sync(); }
+};
+
+/** @brief User exposed API interface to construct grid cooperative group type
+ * object - `grid_group`
+ *
+ * \details User is not allowed to directly construct an object of type
+ * `multi_grid_group`. Instead, he should construct it through this
+ * API function
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+__CG_QUALIFIER__ grid_group this_grid() { return grid_group(internal::grid::size()); }
+
+/** @brief The workgroup (thread-block in CUDA terminology) cooperative group
+ * type
+ *
+ * \details Represents an intra-workgroup cooperative group type where the
+ * participating threads within the group are exactly the same threads
+ * which are participated in the currently executing `workgroup`
+ * @note This is implemented on Linux, under developement
+ * on Windows.
+ */
+class thread_block : public thread_group {
+ // Only these friend functions are allowed to construct an object of thi
+ // class and access its resources
+ friend __CG_QUALIFIER__ thread_block this_thread_block();
+ friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
+ unsigned int tile_size);
+ friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_block& parent,
+ unsigned int tile_size);
+ protected:
+ // Construct a workgroup thread group (through the API this_thread_block())
+ explicit __CG_QUALIFIER__ thread_block(uint32_t size)
+ : thread_group(internal::cg_workgroup, size) {}
+
+ __CG_QUALIFIER__ thread_group new_tiled_group(unsigned int tile_size) const {
+ const bool pow2 = ((tile_size & (tile_size - 1)) == 0);
+ // Invalid tile size, assert
+ if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) {
+ __hip_assert(false && "invalid tile size");
+ }
+
+ thread_group tiledGroup = thread_group(internal::cg_tiled_group, tile_size);
+ tiledGroup.coalesced_info.tiled_info.size = tile_size;
+ tiledGroup.coalesced_info.tiled_info.is_tiled = true;
+ tiledGroup.coalesced_info.tiled_info.meta_group_rank = thread_rank() / tile_size;
+ tiledGroup.coalesced_info.tiled_info.meta_group_size = (size() + tile_size - 1) / tile_size;
+ return tiledGroup;
+ }
+
+ public:
+ // 3-dimensional block index within the grid
+ __CG_STATIC_QUALIFIER__ dim3 group_index() { return internal::workgroup::group_index(); }
+ // 3-dimensional thread index within the block
+ __CG_STATIC_QUALIFIER__ dim3 thread_index() { return internal::workgroup::thread_index(); }
+ __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { return internal::workgroup::thread_rank(); }
+ __CG_STATIC_QUALIFIER__ uint32_t size() { return internal::workgroup::size(); }
+ __CG_STATIC_QUALIFIER__ bool is_valid() { return internal::workgroup::is_valid(); }
+ __CG_STATIC_QUALIFIER__ void sync() { internal::workgroup::sync(); }
+ __CG_QUALIFIER__ dim3 group_dim() { return internal::workgroup::block_dim(); }
+};
+
+/** \brief User exposed API interface to construct workgroup cooperative
+ * group type object - `thread_block`.
+ *
+ * \details User is not allowed to directly construct an object of type
+ * `thread_block`. Instead, he should construct it through this API
+ * function.
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+__CG_QUALIFIER__ thread_block this_thread_block() {
+ return thread_block(internal::workgroup::size());
+}
+
+/** \brief The tiled_group cooperative group type
+ *
+ * \details Represents one tiled thread group in a wavefront.
+ * This group type also supports sub-wave level intrinsics.
+ * @note This is implemented on Linux, under developement
+ * on Windows.
+ */
+
+class tiled_group : public thread_group {
+ private:
+ friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent,
+ unsigned int tile_size);
+ friend __CG_QUALIFIER__ tiled_group tiled_partition(const tiled_group& parent,
+ unsigned int tile_size);
+
+ __CG_QUALIFIER__ tiled_group new_tiled_group(unsigned int tile_size) const {
+ const bool pow2 = ((tile_size & (tile_size - 1)) == 0);
+
+ if (!tile_size || (tile_size > __AMDGCN_WAVEFRONT_SIZE) || !pow2) {
+ __hip_assert(false && "invalid tile size");
+ }
+
+ if (size() <= tile_size) {
+ return *this;
+ }
+
+ tiled_group tiledGroup = tiled_group(tile_size);
+ tiledGroup.coalesced_info.tiled_info.is_tiled = true;
+ return tiledGroup;
+ }
+
+ protected:
+ explicit __CG_QUALIFIER__ tiled_group(unsigned int tileSize)
+ : thread_group(internal::cg_tiled_group, tileSize) {
+ coalesced_info.tiled_info.size = tileSize;
+ coalesced_info.tiled_info.is_tiled = true;
+ }
+
+ public:
+ __CG_QUALIFIER__ unsigned int size() const { return (coalesced_info.tiled_info.size); }
+
+ __CG_QUALIFIER__ unsigned int thread_rank() const {
+ return (internal::workgroup::thread_rank() & (coalesced_info.tiled_info.size - 1));
+ }
+
+ __CG_QUALIFIER__ void sync() const {
+ internal::tiled_group::sync();
+ }
+};
+
+/** \brief The coalesced_group cooperative group type
+ *
+ * \details Represents a active thread group in a wavefront.
+ * This group type also supports sub-wave level intrinsics.
+ * @note This is implemented on Linux, under developement
+ * on Windows.
+ */
+class coalesced_group : public thread_group {
+ private:
+ friend __CG_QUALIFIER__ coalesced_group coalesced_threads();
+ friend __CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, unsigned int tile_size);
+ friend __CG_QUALIFIER__ coalesced_group tiled_partition(const coalesced_group& parent, unsigned int tile_size);
+
+ __CG_QUALIFIER__ coalesced_group new_tiled_group(unsigned int tile_size) const {
+ const bool pow2 = ((tile_size & (tile_size - 1)) == 0);
+
+ if (!tile_size || (tile_size > size()) || !pow2) {
+ return coalesced_group(0);
+ }
+
+ // If a tiled group is passed to be partitioned further into a coalesced_group.
+ // prepare a mask for further partitioning it so that it stays coalesced.
+ if (coalesced_info.tiled_info.is_tiled) {
+ unsigned int base_offset = (thread_rank() & (~(tile_size - 1)));
+ unsigned int masklength = min(static_cast(size()) - base_offset, tile_size);
+ lane_mask member_mask = static_cast(-1) >> (__AMDGCN_WAVEFRONT_SIZE - masklength);
+
+ member_mask <<= (__lane_id() & ~(tile_size - 1));
+ coalesced_group coalesced_tile = coalesced_group(member_mask);
+ coalesced_tile.coalesced_info.tiled_info.is_tiled = true;
+ coalesced_tile.coalesced_info.tiled_info.meta_group_rank = thread_rank() / tile_size;
+ coalesced_tile.coalesced_info.tiled_info.meta_group_size = size() / tile_size;
+ return coalesced_tile;
+ }
+ // Here the parent coalesced_group is not partitioned.
+ else {
+ lane_mask member_mask = 0;
+ unsigned int tile_rank = 0;
+ int lanes_to_skip = ((thread_rank()) / tile_size) * tile_size;
+
+ for (unsigned int i = 0; i < __AMDGCN_WAVEFRONT_SIZE; i++) {
+ lane_mask active = coalesced_info.member_mask & (1 << i);
+ // Make sure the lane is active
+ if (active) {
+ if (lanes_to_skip <= 0 && tile_rank < tile_size) {
+ // Prepare a member_mask that is appropriate for a tile
+ member_mask |= active;
+ tile_rank++;
+ }
+ lanes_to_skip--;
+ }
+ }
+ coalesced_group coalesced_tile = coalesced_group(member_mask);
+ coalesced_tile.coalesced_info.tiled_info.meta_group_rank = thread_rank() / tile_size;
+ coalesced_tile.coalesced_info.tiled_info.meta_group_size =
+ (size() + tile_size - 1) / tile_size;
+ return coalesced_tile;
+ }
+ return coalesced_group(0);
+ }
+
+ protected:
+ // Constructor
+ explicit __CG_QUALIFIER__ coalesced_group(lane_mask member_mask)
+ : thread_group(internal::cg_coalesced_group) {
+ coalesced_info.member_mask = member_mask; // Which threads are active
+ coalesced_info.size = __popcll(coalesced_info.member_mask); // How many threads are active
+ coalesced_info.tiled_info.is_tiled = false; // Not a partitioned group
+ coalesced_info.tiled_info.meta_group_rank = 0;
+ coalesced_info.tiled_info.meta_group_size = 1;
+ }
+
+ public:
+ __CG_QUALIFIER__ unsigned int size() const {
+ return coalesced_info.size;
+ }
+
+ __CG_QUALIFIER__ unsigned int thread_rank() const {
+ return internal::coalesced_group::masked_bit_count(coalesced_info.member_mask);
+ }
+
+ __CG_QUALIFIER__ void sync() const {
+ internal::coalesced_group::sync();
+ }
+
+ __CG_QUALIFIER__ unsigned int meta_group_rank() const {
+ return coalesced_info.tiled_info.meta_group_rank;
+ }
+
+ __CG_QUALIFIER__ unsigned int meta_group_size() const {
+ return coalesced_info.tiled_info.meta_group_size;
+ }
+
+ template
+ __CG_QUALIFIER__ T shfl(T var, int srcRank) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+
+ srcRank = srcRank % static_cast(size());
+
+ int lane = (size() == __AMDGCN_WAVEFRONT_SIZE) ? srcRank
+ : (__AMDGCN_WAVEFRONT_SIZE == 64) ? __fns64(coalesced_info.member_mask, 0, (srcRank + 1))
+ : __fns32(coalesced_info.member_mask, 0, (srcRank + 1));
+
+ return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE);
+ }
+
+ template
+ __CG_QUALIFIER__ T shfl_down(T var, unsigned int lane_delta) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+
+ // Note: The cuda implementation appears to use the remainder of lane_delta
+ // and WARP_SIZE as the shift value rather than lane_delta itself.
+ // This is not described in the documentation and is not done here.
+
+ if (size() == __AMDGCN_WAVEFRONT_SIZE) {
+ return __shfl_down(var, lane_delta, __AMDGCN_WAVEFRONT_SIZE);
+ }
+
+ int lane;
+ if (__AMDGCN_WAVEFRONT_SIZE == 64) {
+ lane = __fns64(coalesced_info.member_mask, __lane_id(), lane_delta + 1);
+ }
+ else {
+ lane = __fns32(coalesced_info.member_mask, __lane_id(), lane_delta + 1);
+ }
+
+ if (lane == -1) {
+ lane = __lane_id();
+ }
+
+ return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE);
+ }
+
+ template
+ __CG_QUALIFIER__ T shfl_up(T var, unsigned int lane_delta) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+
+ // Note: The cuda implementation appears to use the remainder of lane_delta
+ // and WARP_SIZE as the shift value rather than lane_delta itself.
+ // This is not described in the documentation and is not done here.
+
+ if (size() == __AMDGCN_WAVEFRONT_SIZE) {
+ return __shfl_up(var, lane_delta, __AMDGCN_WAVEFRONT_SIZE);
+ }
+
+ int lane;
+ if (__AMDGCN_WAVEFRONT_SIZE == 64) {
+ lane = __fns64(coalesced_info.member_mask, __lane_id(), -(lane_delta + 1));
+ }
+ else if (__AMDGCN_WAVEFRONT_SIZE == 32) {
+ lane = __fns32(coalesced_info.member_mask, __lane_id(), -(lane_delta + 1));
+ }
+
+ if (lane == -1) {
+ lane = __lane_id();
+ }
+
+ return __shfl(var, lane, __AMDGCN_WAVEFRONT_SIZE);
+ }
+};
+
+/** \brief User exposed API to create coalesced groups.
+ *
+ * \details A collective operation that groups all active lanes into a new thread group.
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+
+__CG_QUALIFIER__ coalesced_group coalesced_threads() {
+ return cooperative_groups::coalesced_group(__builtin_amdgcn_read_exec());
+}
+
+/**
+ * Implemenation of all publicly exposed base class APIs
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+__CG_QUALIFIER__ uint32_t thread_group::thread_rank() const {
+ switch (this->_type) {
+ case internal::cg_multi_grid: {
+ return (static_cast(this)->thread_rank());
+ }
+ case internal::cg_grid: {
+ return (static_cast(this)->thread_rank());
+ }
+ case internal::cg_workgroup: {
+ return (static_cast(this)->thread_rank());
+ }
+ case internal::cg_tiled_group: {
+ return (static_cast(this)->thread_rank());
+ }
+ case internal::cg_coalesced_group: {
+ return (static_cast(this)->thread_rank());
+ }
+ default: {
+ __hip_assert(false && "invalid cooperative group type");
+ return -1;
+ }
+ }
+}
+/**
+ * Implemenation of all publicly exposed thread group API
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+__CG_QUALIFIER__ bool thread_group::is_valid() const {
+ switch (this->_type) {
+ case internal::cg_multi_grid: {
+ return (static_cast(this)->is_valid());
+ }
+ case internal::cg_grid: {
+ return (static_cast(this)->is_valid());
+ }
+ case internal::cg_workgroup: {
+ return (static_cast(this)->is_valid());
+ }
+ case internal::cg_tiled_group: {
+ return (static_cast(this)->is_valid());
+ }
+ case internal::cg_coalesced_group: {
+ return (static_cast(this)->is_valid());
+ }
+ default: {
+ __hip_assert(false && "invalid cooperative group type");
+ return false;
+ }
+ }
+}
+/**
+ * Implemenation of all publicly exposed thread group sync API
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+__CG_QUALIFIER__ void thread_group::sync() const {
+ switch (this->_type) {
+ case internal::cg_multi_grid: {
+ static_cast(this)->sync();
+ break;
+ }
+ case internal::cg_grid: {
+ static_cast(this)->sync();
+ break;
+ }
+ case internal::cg_workgroup: {
+ static_cast(this)->sync();
+ break;
+ }
+ case internal::cg_tiled_group: {
+ static_cast(this)->sync();
+ break;
+ }
+ case internal::cg_coalesced_group: {
+ static_cast(this)->sync();
+ break;
+ }
+ default: {
+ __hip_assert(false && "invalid cooperative group type");
+ }
+ }
+}
+
+/**
+ * Implemenation of publicly exposed `wrapper` API on top of basic cooperative
+ * group type APIs
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+template __CG_QUALIFIER__ uint32_t group_size(CGTy const& g) { return g.size(); }
+/**
+ * Implemenation of publicly exposed `wrapper` API on top of basic cooperative
+ * group type APIs
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+template __CG_QUALIFIER__ uint32_t thread_rank(CGTy const& g) {
+ return g.thread_rank();
+}
+/**
+ * Implemenation of publicly exposed `wrapper` API on top of basic cooperative
+ * group type APIs
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+template __CG_QUALIFIER__ bool is_valid(CGTy const& g) { return g.is_valid(); }
+/**
+ * Implemenation of publicly exposed `wrapper` API on top of basic cooperative
+ * group type APIs
+ * @note This function is implemented on Linux, under developement
+ * on Windows.
+ */
+template __CG_QUALIFIER__ void sync(CGTy const& g) { g.sync(); }
+/**
+ * template class tile_base
+ * @note This class is implemented on Linux, under developement
+ * on Windows.
+ */
+template class tile_base {
+ protected:
+ _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize;
+
+ public:
+ // Rank of the thread within this tile
+ _CG_STATIC_CONST_DECL_ unsigned int thread_rank() {
+ return (internal::workgroup::thread_rank() & (numThreads - 1));
+ }
+
+ // Number of threads within this tile
+ __CG_STATIC_QUALIFIER__ unsigned int size() { return numThreads; }
+};
+/**
+ * template class thread_block_tile_base
+ * @note This class is implemented on Linux, under developement
+ * on Windows.
+ */
+template class thread_block_tile_base : public tile_base {
+ static_assert(is_valid_tile_size::value,
+ "Tile size is either not a power of 2 or greater than the wavefront size");
+ using tile_base::numThreads;
+
+ public:
+ __CG_STATIC_QUALIFIER__ void sync() {
+ internal::tiled_group::sync();
+ }
+
+ template __CG_QUALIFIER__ T shfl(T var, int srcRank) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+ return (__shfl(var, srcRank, numThreads));
+ }
+
+ template __CG_QUALIFIER__ T shfl_down(T var, unsigned int lane_delta) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+ return (__shfl_down(var, lane_delta, numThreads));
+ }
+
+ template __CG_QUALIFIER__ T shfl_up(T var, unsigned int lane_delta) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+ return (__shfl_up(var, lane_delta, numThreads));
+ }
+
+ template __CG_QUALIFIER__ T shfl_xor(T var, unsigned int laneMask) const {
+ static_assert(is_valid_type::value, "Neither an integer or float type.");
+ return (__shfl_xor(var, laneMask, numThreads));
+ }
+};
+/** \brief User exposed API that captures the state of the parent group pre-partition
+ */
+template
+class parent_group_info {
+public:
+ // Returns the linear rank of the group within the set of tiles partitioned
+ // from a parent group (bounded by meta_group_size)
+ __CG_STATIC_QUALIFIER__ unsigned int meta_group_rank() {
+ return ParentCGTy::thread_rank() / tileSize;
+ }
+
+ // Returns the number of groups created when the parent group was partitioned.
+ __CG_STATIC_QUALIFIER__ unsigned int meta_group_size() {
+ return (ParentCGTy::size() + tileSize - 1) / tileSize;
+ }
+};
+
+/** \brief Group type - thread_block_tile
+ *
+ * \details Represents one tile of thread group.
+ * @note This type is implemented on Linux, under developement
+ * on Windows.
+ */
+template
+class thread_block_tile_type : public thread_block_tile_base,
+ public tiled_group,
+ public parent_group_info {
+ _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize;
+ typedef thread_block_tile_base tbtBase;
+ protected:
+ __CG_QUALIFIER__ thread_block_tile_type() : tiled_group(numThreads) {
+ coalesced_info.tiled_info.size = numThreads;
+ coalesced_info.tiled_info.is_tiled = true;
+ }
+ public:
+ using tbtBase::size;
+ using tbtBase::sync;
+ using tbtBase::thread_rank;
+};
+
+// Partial template specialization
+template
+class thread_block_tile_type : public thread_block_tile_base,
+ public tiled_group
+ {
+ _CG_STATIC_CONST_DECL_ unsigned int numThreads = tileSize;
+
+ typedef thread_block_tile_base tbtBase;
+
+ protected:
+
+ __CG_QUALIFIER__ thread_block_tile_type(unsigned int meta_group_rank, unsigned int meta_group_size)
+ : tiled_group(numThreads) {
+ coalesced_info.tiled_info.size = numThreads;
+ coalesced_info.tiled_info.is_tiled = true;
+ coalesced_info.tiled_info.meta_group_rank = meta_group_rank;
+ coalesced_info.tiled_info.meta_group_size = meta_group_size;
+ }
+
+ public:
+ using tbtBase::size;
+ using tbtBase::sync;
+ using tbtBase::thread_rank;
+
+ __CG_QUALIFIER__ unsigned int meta_group_rank() const {
+ return coalesced_info.tiled_info.meta_group_rank;
+ }
+
+ __CG_QUALIFIER__ unsigned int meta_group_size() const {
+ return coalesced_info.tiled_info.meta_group_size;
+ }
+// end of operative group
+/**
+* @}
+*/
+};
+
+
+/** \brief User exposed API to partition groups.
+ *
+ * \details A collective operation that partitions the parent group into a one-dimensional,
+ * row-major, tiling of subgroups.
+ */
+
+__CG_QUALIFIER__ thread_group tiled_partition(const thread_group& parent, unsigned int tile_size) {
+ if (parent.cg_type() == internal::cg_tiled_group) {
+ const tiled_group* cg = static_cast(&parent);
+ return cg->new_tiled_group(tile_size);
+ }
+ else if(parent.cg_type() == internal::cg_coalesced_group) {
+ const coalesced_group* cg = static_cast(&parent);
+ return cg->new_tiled_group(tile_size);
+ }
+ else {
+ const thread_block* tb = static_cast(&parent);
+ return tb->new_tiled_group(tile_size);
+ }
+}
+
+// Thread block type overload
+__CG_QUALIFIER__ thread_group tiled_partition(const thread_block& parent, unsigned int tile_size) {
+ return (parent.new_tiled_group(tile_size));
+}
+
+__CG_QUALIFIER__ tiled_group tiled_partition(const tiled_group& parent, unsigned int tile_size) {
+ return (parent.new_tiled_group(tile_size));
+}
+
+// If a coalesced group is passed to be partitioned, it should remain coalesced
+__CG_QUALIFIER__ coalesced_group tiled_partition(const coalesced_group& parent, unsigned int tile_size) {
+ return (parent.new_tiled_group(tile_size));
+}
+
+template class thread_block_tile;
+
+namespace impl {
+template class thread_block_tile_internal;
+
+template
+class thread_block_tile_internal : public thread_block_tile_type {
+ protected:
+ template
+ __CG_QUALIFIER__ thread_block_tile_internal(
+ const thread_block_tile_internal& g)
+ : thread_block_tile_type(g.meta_group_rank(), g.meta_group_size()) {}
+
+ __CG_QUALIFIER__ thread_block_tile_internal(const thread_block& g)
+ : thread_block_tile_type() {}
+};
+} // namespace impl
+
+template
+class thread_block_tile : public impl::thread_block_tile_internal {
+ protected:
+ __CG_QUALIFIER__ thread_block_tile(const ParentCGTy& g)
+ : impl::thread_block_tile_internal(g) {}
+
+ public:
+ __CG_QUALIFIER__ operator thread_block_tile() const {
+ return thread_block_tile(*this);
+ }
+};
+
+
+template
+class thread_block_tile : public impl::thread_block_tile_internal {
+ template friend class thread_block_tile;
+
+ protected:
+ public:
+ template
+ __CG_QUALIFIER__ thread_block_tile(const thread_block_tile& g)
+ : impl::thread_block_tile_internal(g) {}
+};
+
+template class thread_block_tile;
+
+namespace impl {
+template struct tiled_partition_internal;
+
+template
+struct tiled_partition_internal : public thread_block_tile {
+ __CG_QUALIFIER__ tiled_partition_internal(const thread_block& g)
+ : thread_block_tile(g) {}
+};
+
+} // namespace impl
+
+/** \brief User exposed API to partition groups.
+ *
+ * \details This constructs a templated class derieved from thread_group.
+ * The template defines tile size of the new thread group at compile time.
+ */
+template
+__CG_QUALIFIER__ thread_block_tile tiled_partition(const ParentCGTy& g) {
+ static_assert(is_valid_tile_size::value,
+ "Tiled partition with size > wavefront size. Currently not supported ");
+ return impl::tiled_partition_internal(g);
+}
+} // namespace cooperative_groups
+
+#endif // __cplusplus
+#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_COOPERATIVE_GROUPS_H
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_fp16.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_fp16.h
new file mode 100644
index 0000000000000000000000000000000000000000..62d88a3756b1ac2607593ca7c52a44a5034db6a2
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_fp16.h
@@ -0,0 +1,1809 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#pragma once
+#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H
+#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H
+
+#if defined(__HIPCC_RTC__)
+ #define __HOST_DEVICE__ __device__
+#else
+ #define __HOST_DEVICE__ __host__ __device__
+ #include
+ #include "hip/amd_detail/host_defines.h"
+ #include
+ #if defined(__cplusplus)
+ #include
+ #include
+ #include
+#endif
+#endif // !defined(__HIPCC_RTC__)
+
+#if defined(__clang__) && defined(__HIP__)
+ typedef _Float16 _Float16_2 __attribute__((ext_vector_type(2)));
+
+ struct __half_raw {
+ union {
+ static_assert(sizeof(_Float16) == sizeof(unsigned short), "");
+
+ _Float16 data;
+ unsigned short x;
+ };
+ };
+
+ struct __half2_raw {
+ union {
+ static_assert(sizeof(_Float16_2) == sizeof(unsigned short[2]), "");
+
+ struct {
+ __half_raw x;
+ __half_raw y;
+ };
+ _Float16_2 data;
+ };
+ };
+
+ #if defined(__cplusplus)
+ #if !defined(__HIPCC_RTC__)
+ #include "hip_fp16_math_fwd.h"
+ #include "amd_hip_vector_types.h"
+ #include "host_defines.h"
+ #include "amd_device_functions.h"
+ #include "amd_warp_functions.h"
+ #endif
+ namespace std
+ {
+ template<> struct is_floating_point<_Float16> : std::true_type {};
+ }
+
+ template
+ using Enable_if_t = typename std::enable_if::type;
+
+ // BEGIN STRUCT __HALF
+ struct __half {
+ protected:
+ union {
+ static_assert(sizeof(_Float16) == sizeof(unsigned short), "");
+
+ _Float16 data;
+ unsigned short __x;
+ };
+ public:
+ // CREATORS
+ __HOST_DEVICE__
+ __half() = default;
+ __HOST_DEVICE__
+ __half(const __half_raw& x) : data{x.data} {}
+ #if !defined(__HIP_NO_HALF_CONVERSIONS__)
+ __HOST_DEVICE__
+ __half(decltype(data) x) : data{x} {}
+ template<
+ typename T,
+ Enable_if_t{}>* = nullptr>
+ __HOST_DEVICE__
+ __half(T x) : data{static_cast<_Float16>(x)} {}
+ #endif
+ __HOST_DEVICE__
+ __half(const __half&) = default;
+ __HOST_DEVICE__
+ __half(__half&&) = default;
+ __HOST_DEVICE__
+ ~__half() = default;
+
+ // CREATORS - DEVICE ONLY
+ #if !defined(__HIP_NO_HALF_CONVERSIONS__)
+ template<
+ typename T, Enable_if_t{}>* = nullptr>
+ __HOST_DEVICE__
+ __half(T x) : data{static_cast<_Float16>(x)} {}
+ #endif
+
+ // MANIPULATORS
+ __HOST_DEVICE__
+ __half& operator=(const __half&) = default;
+ __HOST_DEVICE__
+ __half& operator=(__half&&) = default;
+ __HOST_DEVICE__
+ __half& operator=(const __half_raw& x)
+ {
+ data = x.data;
+ return *this;
+ }
+ __HOST_DEVICE__
+ volatile __half& operator=(const __half_raw& x) volatile
+ {
+ data = x.data;
+ return *this;
+ }
+ volatile __half& operator=(const volatile __half_raw& x) volatile
+ {
+ data = x.data;
+ return *this;
+ }
+ __half& operator=(__half_raw&& x)
+ {
+ data = x.data;
+ return *this;
+ }
+ volatile __half& operator=(__half_raw&& x) volatile
+ {
+ data = x.data;
+ return *this;
+ }
+ volatile __half& operator=(volatile __half_raw&& x) volatile
+ {
+ data = x.data;
+ return *this;
+ }
+ #if !defined(__HIP_NO_HALF_CONVERSIONS__)
+ template<
+ typename T,
+ Enable_if_t{}>* = nullptr>
+ __HOST_DEVICE__
+ __half& operator=(T x)
+ {
+ data = static_cast<_Float16>(x);
+ return *this;
+ }
+ #endif
+
+ // MANIPULATORS - DEVICE ONLY
+ #if !defined(__HIP_NO_HALF_CONVERSIONS__)
+ template<
+ typename T, Enable_if_t{}>* = nullptr>
+ __device__
+ __half& operator=(T x)
+ {
+ data = static_cast<_Float16>(x);
+ return *this;
+ }
+ #endif
+
+ #if !defined(__HIP_NO_HALF_OPERATORS__)
+ __device__
+ __half& operator+=(const __half& x)
+ {
+ data += x.data;
+ return *this;
+ }
+ __device__
+ __half& operator-=(const __half& x)
+ {
+ data -= x.data;
+ return *this;
+ }
+ __device__
+ __half& operator*=(const __half& x)
+ {
+ data *= x.data;
+ return *this;
+ }
+ __device__
+ __half& operator/=(const __half& x)
+ {
+ data /= x.data;
+ return *this;
+ }
+ __device__
+ __half& operator++() { ++data; return *this; }
+ __device__
+ __half operator++(int)
+ {
+ __half tmp{*this};
+ ++*this;
+ return tmp;
+ }
+ __device__
+ __half& operator--() { --data; return *this; }
+ __device__
+ __half operator--(int)
+ {
+ __half tmp{*this};
+ --*this;
+ return tmp;
+ }
+ #endif
+
+ // ACCESSORS
+ #if !defined(__HIP_NO_HALF_CONVERSIONS__)
+ template<
+ typename T,
+ Enable_if_t{}>* = nullptr>
+ __HOST_DEVICE__
+ operator T() const { return data; }
+ #endif
+ __HOST_DEVICE__
+ operator __half_raw() const { return __half_raw{data}; }
+ __HOST_DEVICE__
+ operator __half_raw() const volatile
+ {
+ return __half_raw{data};
+ }
+
+ #if !defined(__HIP_NO_HALF_CONVERSIONS__)
+ template<
+ typename T, Enable_if_t{}>* = nullptr>
+ __HOST_DEVICE__
+ operator T() const { return data; }
+ #endif
+
+ #if !defined(__HIP_NO_HALF_OPERATORS__)
+ __device__
+ __half operator+() const { return *this; }
+ __device__
+ __half operator-() const
+ {
+ __half tmp{*this};
+ tmp.data = -tmp.data;
+ return tmp;
+ }
+ #endif
+
+ // FRIENDS
+ #if !defined(__HIP_NO_HALF_OPERATORS__)
+ friend
+ inline
+ __device__
+ __half operator+(const __half& x, const __half& y)
+ {
+ return __half{x} += y;
+ }
+ friend
+ inline
+ __device__
+ __half operator-(const __half& x, const __half& y)
+ {
+ return __half{x} -= y;
+ }
+ friend
+ inline
+ __device__
+ __half operator*(const __half& x, const __half& y)
+ {
+ return __half{x} *= y;
+ }
+ friend
+ inline
+ __device__
+ __half operator/(const __half& x, const __half& y)
+ {
+ return __half{x} /= y;
+ }
+ friend
+ inline
+ __device__
+ bool operator==(const __half& x, const __half& y)
+ {
+ return x.data == y.data;
+ }
+ friend
+ inline
+ __device__
+ bool operator!=(const __half& x, const __half& y)
+ {
+ return !(x == y);
+ }
+ friend
+ inline
+ __device__
+ bool operator<(const __half& x, const __half& y)
+ {
+ return x.data < y.data;
+ }
+ friend
+ inline
+ __device__
+ bool operator>(const __half& x, const __half& y)
+ {
+ return y.data < x.data;
+ }
+ friend
+ inline
+ __device__
+ bool operator<=(const __half& x, const __half& y)
+ {
+ return !(y < x);
+ }
+ friend
+ inline
+ __device__
+ bool operator>=(const __half& x, const __half& y)
+ {
+ return !(x < y);
+ }
+ #endif // !defined(__HIP_NO_HALF_OPERATORS__)
+ };
+ // END STRUCT __HALF
+
+ // BEGIN STRUCT __HALF2
+ struct __half2 {
+ public:
+ union {
+ static_assert(
+ sizeof(_Float16_2) == sizeof(unsigned short[2]), "");
+
+ struct {
+ __half x;
+ __half y;
+ };
+ _Float16_2 data;
+ };
+
+ // CREATORS
+ __HOST_DEVICE__
+ __half2() = default;
+ __HOST_DEVICE__
+ __half2(const __half2_raw& xx) : data{xx.data} {}
+ __HOST_DEVICE__
+ __half2(decltype(data) xx) : data{xx} {}
+ __HOST_DEVICE__
+ __half2(const __half& xx, const __half& yy)
+ :
+ data{static_cast<__half_raw>(xx).data,
+ static_cast<__half_raw>(yy).data}
+ {}
+ __HOST_DEVICE__
+ __half2(const __half2&) = default;
+ __HOST_DEVICE__
+ __half2(__half2&&) = default;
+ __HOST_DEVICE__
+ ~__half2() = default;
+
+ // MANIPULATORS
+ __HOST_DEVICE__
+ __half2& operator=(const __half2&) = default;
+ __HOST_DEVICE__
+ __half2& operator=(__half2&&) = default;
+ __HOST_DEVICE__
+ __half2& operator=(const __half2_raw& xx)
+ {
+ data = xx.data;
+ return *this;
+ }
+
+ // MANIPULATORS - DEVICE ONLY
+ #if !defined(__HIP_NO_HALF_OPERATORS__)
+ __device__
+ __half2& operator+=(const __half2& xx)
+ {
+ data += xx.data;
+ return *this;
+ }
+ __device__
+ __half2& operator-=(const __half2& xx)
+ {
+ data -= xx.data;
+ return *this;
+ }
+ __device__
+ __half2& operator*=(const __half2& xx)
+ {
+ data *= xx.data;
+ return *this;
+ }
+ __device__
+ __half2& operator/=(const __half2& xx)
+ {
+ data /= xx.data;
+ return *this;
+ }
+ __device__
+ __half2& operator++() { return *this += _Float16_2{1, 1}; }
+ __device__
+ __half2 operator++(int)
+ {
+ __half2 tmp{*this};
+ ++*this;
+ return tmp;
+ }
+ __device__
+ __half2& operator--() { return *this -= _Float16_2{1, 1}; }
+ __device__
+ __half2 operator--(int)
+ {
+ __half2 tmp{*this};
+ --*this;
+ return tmp;
+ }
+ #endif
+
+ // ACCESSORS
+ __HOST_DEVICE__
+ operator decltype(data)() const { return data; }
+ __HOST_DEVICE__
+ operator __half2_raw() const {
+ __half2_raw r;
+ r.data = data;
+ return r;
+ }
+
+ // ACCESSORS - DEVICE ONLY
+ #if !defined(__HIP_NO_HALF_OPERATORS__)
+ __device__
+ __half2 operator+() const { return *this; }
+ __device__
+ __half2 operator-() const
+ {
+ __half2 tmp{*this};
+ tmp.data = -tmp.data;
+ return tmp;
+ }
+ #endif
+
+ // FRIENDS
+ #if !defined(__HIP_NO_HALF_OPERATORS__)
+ friend
+ inline
+ __device__
+ __half2 operator+(const __half2& xx, const __half2& yy)
+ {
+ return __half2{xx} += yy;
+ }
+ friend
+ inline
+ __device__
+ __half2 operator-(const __half2& xx, const __half2& yy)
+ {
+ return __half2{xx} -= yy;
+ }
+ friend
+ inline
+ __device__
+ __half2 operator*(const __half2& xx, const __half2& yy)
+ {
+ return __half2{xx} *= yy;
+ }
+ friend
+ inline
+ __device__
+ __half2 operator/(const __half2& xx, const __half2& yy)
+ {
+ return __half2{xx} /= yy;
+ }
+ friend
+ inline
+ __device__
+ bool operator==(const __half2& xx, const __half2& yy)
+ {
+ auto r = xx.data == yy.data;
+ return r.x != 0 && r.y != 0;
+ }
+ friend
+ inline
+ __device__
+ bool operator!=(const __half2& xx, const __half2& yy)
+ {
+ return !(xx == yy);
+ }
+ friend
+ inline
+ __device__
+ bool operator<(const __half2& xx, const __half2& yy)
+ {
+ auto r = xx.data < yy.data;
+ return r.x != 0 && r.y != 0;
+ }
+ friend
+ inline
+ __device__
+ bool operator>(const __half2& xx, const __half2& yy)
+ {
+ return yy < xx;
+ }
+ friend
+ inline
+ __device__
+ bool operator<=(const __half2& xx, const __half2& yy)
+ {
+ return !(yy < xx);
+ }
+ friend
+ inline
+ __device__
+ bool operator>=(const __half2& xx, const __half2& yy)
+ {
+ return !(xx < yy);
+ }
+ #endif // !defined(__HIP_NO_HALF_OPERATORS__)
+ };
+ // END STRUCT __HALF2
+
+ namespace
+ {
+ inline
+ __HOST_DEVICE__
+ __half2 make_half2(__half x, __half y)
+ {
+ return __half2{x, y};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half __low2half(__half2 x)
+ {
+ return __half{__half_raw{static_cast<__half2_raw>(x).data.x}};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half __high2half(__half2 x)
+ {
+ return __half{__half_raw{static_cast<__half2_raw>(x).data.y}};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __half2half2(__half x)
+ {
+ return __half2{x, x};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __halves2half2(__half x, __half y)
+ {
+ return __half2{x, y};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __low2half2(__half2 x)
+ {
+ return __half2{
+ _Float16_2{
+ static_cast<__half2_raw>(x).data.x,
+ static_cast<__half2_raw>(x).data.x}};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __high2half2(__half2 x)
+ {
+ return __half2{
+ _Float16_2{
+ static_cast<__half2_raw>(x).data.y,
+ static_cast<__half2_raw>(x).data.y}};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __lows2half2(__half2 x, __half2 y)
+ {
+ return __half2{
+ _Float16_2{
+ static_cast<__half2_raw>(x).data.x,
+ static_cast<__half2_raw>(y).data.x}};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __highs2half2(__half2 x, __half2 y)
+ {
+ return __half2{
+ _Float16_2{
+ static_cast<__half2_raw>(x).data.y,
+ static_cast<__half2_raw>(y).data.y}};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __lowhigh2highlow(__half2 x)
+ {
+ return __half2{
+ _Float16_2{
+ static_cast<__half2_raw>(x).data.y,
+ static_cast<__half2_raw>(x).data.x}};
+ }
+
+ // Bitcasts
+ inline
+ __device__
+ short __half_as_short(__half x)
+ {
+ return static_cast<__half_raw>(x).x;
+ }
+
+ inline
+ __device__
+ unsigned short __half_as_ushort(__half x)
+ {
+ return static_cast<__half_raw>(x).x;
+ }
+
+ inline
+ __device__
+ __half __short_as_half(short x)
+ {
+ __half_raw r; r.x = x;
+ return r;
+ }
+
+ inline
+ __device__
+ __half __ushort_as_half(unsigned short x)
+ {
+ __half_raw r; r.x = x;
+ return r;
+ }
+
+ // float -> half | half2
+ inline
+ __HOST_DEVICE__
+ __half __float2half(float x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half __float2half_rn(float x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ #if !defined(__HIPCC_RTC__)
+ // TODO: rounding behaviour is not correct for host functions.
+ inline
+ __host__
+ __half __float2half_rz(float x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __host__
+ __half __float2half_rd(float x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __host__
+ __half __float2half_ru(float x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ #endif
+ inline
+ __device__
+ __half __float2half_rz(float x)
+ {
+ return __half_raw{__ocml_cvtrtz_f16_f32(x)};
+ }
+ inline
+ __device__
+ __half __float2half_rd(float x)
+ {
+ return __half_raw{__ocml_cvtrtn_f16_f32(x)};
+ }
+ inline
+ __device__
+ __half __float2half_ru(float x)
+ {
+ return __half_raw{__ocml_cvtrtp_f16_f32(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __float2half2_rn(float x)
+ {
+ return __half2{
+ _Float16_2{
+ static_cast<_Float16>(x), static_cast<_Float16>(x)}};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __floats2half2_rn(float x, float y)
+ {
+ return __half2{_Float16_2{
+ static_cast<_Float16>(x), static_cast<_Float16>(y)}};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __float22half2_rn(float2 x)
+ {
+ return __floats2half2_rn(x.x, x.y);
+ }
+
+ // half | half2 -> float
+ inline
+ __HOST_DEVICE__
+ float __half2float(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __HOST_DEVICE__
+ float __low2float(__half2 x)
+ {
+ return static_cast<__half2_raw>(x).data.x;
+ }
+ inline
+ __HOST_DEVICE__
+ float __high2float(__half2 x)
+ {
+ return static_cast<__half2_raw>(x).data.y;
+ }
+ inline
+ __HOST_DEVICE__
+ float2 __half22float2(__half2 x)
+ {
+ return make_float2(
+ static_cast<__half2_raw>(x).data.x,
+ static_cast<__half2_raw>(x).data.y);
+ }
+
+ // half -> int
+ inline
+ __device__
+ int __half2int_rn(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ int __half2int_rz(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ int __half2int_rd(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ int __half2int_ru(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+
+ // int -> half
+ inline
+ __device__
+ __half __int2half_rn(int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __int2half_rz(int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __int2half_rd(int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __int2half_ru(int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+
+ // half -> short
+ inline
+ __device__
+ short __half2short_rn(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ short __half2short_rz(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ short __half2short_rd(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ short __half2short_ru(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+
+ // short -> half
+ inline
+ __device__
+ __half __short2half_rn(short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __short2half_rz(short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __short2half_rd(short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __short2half_ru(short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+
+ // half -> long long
+ inline
+ __device__
+ long long __half2ll_rn(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ long long __half2ll_rz(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ long long __half2ll_rd(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ long long __half2ll_ru(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+
+ // long long -> half
+ inline
+ __device__
+ __half __ll2half_rn(long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ll2half_rz(long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ll2half_rd(long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ll2half_ru(long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+
+ // half -> unsigned int
+ inline
+ __device__
+ unsigned int __half2uint_rn(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned int __half2uint_rz(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned int __half2uint_rd(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned int __half2uint_ru(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+
+ // unsigned int -> half
+ inline
+ __device__
+ __half __uint2half_rn(unsigned int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __uint2half_rz(unsigned int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __uint2half_rd(unsigned int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __uint2half_ru(unsigned int x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+
+ // half -> unsigned short
+ inline
+ __device__
+ unsigned short __half2ushort_rn(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned short __half2ushort_rz(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned short __half2ushort_rd(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned short __half2ushort_ru(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+
+ // unsigned short -> half
+ inline
+ __device__
+ __half __ushort2half_rn(unsigned short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ushort2half_rz(unsigned short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ushort2half_rd(unsigned short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ushort2half_ru(unsigned short x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+
+ // half -> unsigned long long
+ inline
+ __device__
+ unsigned long long __half2ull_rn(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned long long __half2ull_rz(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned long long __half2ull_rd(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+ inline
+ __device__
+ unsigned long long __half2ull_ru(__half x)
+ {
+ return static_cast<__half_raw>(x).data;
+ }
+
+ // unsigned long long -> half
+ inline
+ __device__
+ __half __ull2half_rn(unsigned long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ull2half_rz(unsigned long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ull2half_rd(unsigned long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+ inline
+ __device__
+ __half __ull2half_ru(unsigned long long x)
+ {
+ return __half_raw{static_cast<_Float16>(x)};
+ }
+
+ // Load primitives
+ inline
+ __device__
+ __half __ldg(const __half* ptr) { return *ptr; }
+ inline
+ __device__
+ __half __ldcg(const __half* ptr) { return *ptr; }
+ inline
+ __device__
+ __half __ldca(const __half* ptr) { return *ptr; }
+ inline
+ __device__
+ __half __ldcs(const __half* ptr) { return *ptr; }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __ldg(const __half2* ptr) { return *ptr; }
+ inline
+ __HOST_DEVICE__
+ __half2 __ldcg(const __half2* ptr) { return *ptr; }
+ inline
+ __HOST_DEVICE__
+ __half2 __ldca(const __half2* ptr) { return *ptr; }
+ inline
+ __HOST_DEVICE__
+ __half2 __ldcs(const __half2* ptr) { return *ptr; }
+
+ // Relations
+ inline
+ __device__
+ bool __heq(__half x, __half y)
+ {
+ return static_cast<__half_raw>(x).data ==
+ static_cast<__half_raw>(y).data;
+ }
+ inline
+ __device__
+ bool __hne(__half x, __half y)
+ {
+ return static_cast<__half_raw>(x).data !=
+ static_cast<__half_raw>(y).data;
+ }
+ inline
+ __device__
+ bool __hle(__half x, __half y)
+ {
+ return static_cast<__half_raw>(x).data <=
+ static_cast<__half_raw>(y).data;
+ }
+ inline
+ __device__
+ bool __hge(__half x, __half y)
+ {
+ return static_cast<__half_raw>(x).data >=
+ static_cast<__half_raw>(y).data;
+ }
+ inline
+ __device__
+ bool __hlt(__half x, __half y)
+ {
+ return static_cast<__half_raw>(x).data <
+ static_cast<__half_raw>(y).data;
+ }
+ inline
+ __device__
+ bool __hgt(__half x, __half y)
+ {
+ return static_cast<__half_raw>(x).data >
+ static_cast<__half_raw>(y).data;
+ }
+ inline __device__
+ bool __hequ(__half x, __half y) {
+ return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data) &&
+ !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data);
+ }
+ inline __device__
+ bool __hneu(__half x, __half y) {
+ return !(static_cast<__half_raw>(x).data == static_cast<__half_raw>(y).data);
+ }
+ inline __device__
+ bool __hleu(__half x, __half y) {
+ return !(static_cast<__half_raw>(x).data > static_cast<__half_raw>(y).data);
+ }
+ inline
+ __device__
+ bool __hgeu(__half x, __half y) {
+ return !(static_cast<__half_raw>(x).data < static_cast<__half_raw>(y).data);
+ }
+ inline
+ __device__
+ bool __hltu(__half x, __half y) {
+ return !(static_cast<__half_raw>(x).data >= static_cast<__half_raw>(y).data);
+ }
+ inline
+ __device__
+ bool __hgtu(__half x, __half y) {
+ return !(static_cast<__half_raw>(x).data <= static_cast<__half_raw>(y).data);
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __heq2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(x).data ==
+ static_cast<__half2_raw>(y).data;
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hne2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(x).data !=
+ static_cast<__half2_raw>(y).data;
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hle2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(x).data <=
+ static_cast<__half2_raw>(y).data;
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hge2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(x).data >=
+ static_cast<__half2_raw>(y).data;
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hlt2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(x).data <
+ static_cast<__half2_raw>(y).data;
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hgt2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(x).data >
+ static_cast<__half2_raw>(y).data;
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline __HOST_DEVICE__
+ __half2 __hequ2(__half2 x, __half2 y) {
+ auto r = !(static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data) &&
+ !(static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data);
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hneu2(__half2 x, __half2 y) {
+ auto r = !(static_cast<__half2_raw>(x).data == static_cast<__half2_raw>(y).data);
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hleu2(__half2 x, __half2 y) {
+ auto r = !(static_cast<__half2_raw>(x).data > static_cast<__half2_raw>(y).data);
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hgeu2(__half2 x, __half2 y) {
+ auto r = !(static_cast<__half2_raw>(x).data < static_cast<__half2_raw>(y).data);
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hltu2(__half2 x, __half2 y) {
+ auto r = !(static_cast<__half2_raw>(x).data >= static_cast<__half2_raw>(y).data);
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hgtu2(__half2 x, __half2 y) {
+ auto r = !(static_cast<__half2_raw>(x).data <= static_cast<__half2_raw>(y).data);
+ return __builtin_convertvector(-r, _Float16_2);
+ }
+
+ inline
+ __HOST_DEVICE__
+ bool __hbeq2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__heq2(x, y));
+ return r.data.x != 0 && r.data.y != 0;
+ }
+ inline
+ __HOST_DEVICE__
+ bool __hbne2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hne2(x, y));
+ return r.data.x != 0 && r.data.y != 0;
+ }
+ inline
+ __HOST_DEVICE__
+ bool __hble2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hle2(x, y));
+ return r.data.x != 0 && r.data.y != 0;
+ }
+ inline
+ __HOST_DEVICE__
+ bool __hbge2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hge2(x, y));
+ return r.data.x != 0 && r.data.y != 0;
+ }
+ inline
+ __HOST_DEVICE__
+ bool __hblt2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hlt2(x, y));
+ return r.data.x != 0 && r.data.y != 0;
+ }
+ inline
+ __HOST_DEVICE__
+ bool __hbgt2(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hgt2(x, y));
+ return r.data.x != 0 && r.data.y != 0;
+ }
+ inline
+ __HOST_DEVICE__
+ bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); }
+ inline
+ __HOST_DEVICE__
+ bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); }
+ inline
+ __HOST_DEVICE__
+ bool __hbleu2(__half2 x, __half2 y) { return __hble2(x, y); }
+ inline
+ __HOST_DEVICE__
+ bool __hbgeu2(__half2 x, __half2 y) { return __hbge2(x, y); }
+ inline
+ __HOST_DEVICE__
+ bool __hbltu2(__half2 x, __half2 y) { return __hblt2(x, y); }
+ inline
+ __HOST_DEVICE__
+ bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); }
+ inline
+ __device__
+ __half __hmax(const __half x, const __half y) {
+ return __half_raw{__ocml_fmax_f16(static_cast<__half_raw>(x).data,
+ static_cast<__half_raw>(y).data)};
+ }
+ inline
+ __device__
+ __half __hmax_nan(const __half x, const __half y) {
+ if(__ocml_isnan_f16(static_cast<__half_raw>(x).data)) {
+ return x;
+ } else if (__ocml_isnan_f16(static_cast<__half_raw>(y).data)) {
+ return y;
+ }
+ return __hmax(x, y);
+ }
+ inline
+ __device__
+ __half __hmin(const __half x, const __half y) {
+ return __half_raw{__ocml_fmin_f16(static_cast<__half_raw>(x).data,
+ static_cast<__half_raw>(y).data)};
+ }
+ inline
+ __device__
+ __half __hmin_nan(const __half x, const __half y) {
+ if(__ocml_isnan_f16(static_cast<__half_raw>(x).data)) {
+ return x;
+ } else if (__ocml_isnan_f16(static_cast<__half_raw>(y).data)) {
+ return y;
+ }
+ return __hmin(x, y);
+ }
+
+ // Arithmetic
+ inline
+ __device__
+ __half __clamp_01(__half x)
+ {
+ auto r = static_cast<__half_raw>(x);
+
+ if (__hlt(x, __half_raw{0})) return __half_raw{0};
+ if (__hlt(__half_raw{1}, x)) return __half_raw{1};
+ return r;
+ }
+
+ inline
+ __device__
+ __half __hadd(__half x, __half y)
+ {
+ return __half_raw{
+ static_cast<__half_raw>(x).data +
+ static_cast<__half_raw>(y).data};
+ }
+ inline
+ __device__
+ __half __habs(__half x)
+ {
+ return __half_raw{
+ __ocml_fabs_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half __hsub(__half x, __half y)
+ {
+ return __half_raw{
+ static_cast<__half_raw>(x).data -
+ static_cast<__half_raw>(y).data};
+ }
+ inline
+ __device__
+ __half __hmul(__half x, __half y)
+ {
+ return __half_raw{
+ static_cast<__half_raw>(x).data *
+ static_cast<__half_raw>(y).data};
+ }
+ inline
+ __device__
+ __half __hadd_sat(__half x, __half y)
+ {
+ return __clamp_01(__hadd(x, y));
+ }
+ inline
+ __device__
+ __half __hsub_sat(__half x, __half y)
+ {
+ return __clamp_01(__hsub(x, y));
+ }
+ inline
+ __device__
+ __half __hmul_sat(__half x, __half y)
+ {
+ return __clamp_01(__hmul(x, y));
+ }
+ inline
+ __device__
+ __half __hfma(__half x, __half y, __half z)
+ {
+ return __half_raw{__ocml_fma_f16(
+ static_cast<__half_raw>(x).data,
+ static_cast<__half_raw>(y).data,
+ static_cast<__half_raw>(z).data)};
+ }
+ inline
+ __device__
+ __half __hfma_sat(__half x, __half y, __half z)
+ {
+ return __clamp_01(__hfma(x, y, z));
+ }
+ inline
+ __device__
+ __half __hdiv(__half x, __half y)
+ {
+ return __half_raw{
+ static_cast<__half_raw>(x).data /
+ static_cast<__half_raw>(y).data};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 __hadd2(__half2 x, __half2 y)
+ {
+ return __half2{
+ static_cast<__half2_raw>(x).data +
+ static_cast<__half2_raw>(y).data};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __habs2(__half2 x)
+ {
+ return __half2{
+ __ocml_fabs_2f16(static_cast<__half2_raw>(x).data)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hsub2(__half2 x, __half2 y)
+ {
+ return __half2{
+ static_cast<__half2_raw>(x).data -
+ static_cast<__half2_raw>(y).data};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hmul2(__half2 x, __half2 y)
+ {
+ return __half2{
+ static_cast<__half2_raw>(x).data *
+ static_cast<__half2_raw>(y).data};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hadd2_sat(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hadd2(x, y));
+ return __half2{
+ __clamp_01(__half_raw{r.data.x}),
+ __clamp_01(__half_raw{r.data.y})};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hsub2_sat(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hsub2(x, y));
+ return __half2{
+ __clamp_01(__half_raw{r.data.x}),
+ __clamp_01(__half_raw{r.data.y})};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hmul2_sat(__half2 x, __half2 y)
+ {
+ auto r = static_cast<__half2_raw>(__hmul2(x, y));
+ return __half2{
+ __clamp_01(__half_raw{r.data.x}),
+ __clamp_01(__half_raw{r.data.y})};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hfma2(__half2 x, __half2 y, __half2 z)
+ {
+ return __half2{__ocml_fma_2f16(x, y, z)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z)
+ {
+ auto r = static_cast<__half2_raw>(__hfma2(x, y, z));
+ return __half2{
+ __clamp_01(__half_raw{r.data.x}),
+ __clamp_01(__half_raw{r.data.y})};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __h2div(__half2 x, __half2 y)
+ {
+ return __half2{
+ static_cast<__half2_raw>(x).data /
+ static_cast<__half2_raw>(y).data};
+ }
+
+ // Math functions
+ #if defined(__clang__) && defined(__HIP__)
+ inline
+ __device__
+ float amd_mixed_dot(__half2 a, __half2 b, float c, bool saturate) {
+ return __ockl_fdot2(static_cast<__half2_raw>(a).data,
+ static_cast<__half2_raw>(b).data,
+ c, saturate);
+ }
+ #endif
+ inline
+ __device__
+ __half htrunc(__half x)
+ {
+ return __half_raw{
+ __ocml_trunc_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hceil(__half x)
+ {
+ return __half_raw{
+ __ocml_ceil_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hfloor(__half x)
+ {
+ return __half_raw{
+ __ocml_floor_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hrint(__half x)
+ {
+ return __half_raw{
+ __ocml_rint_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hsin(__half x)
+ {
+ return __half_raw{
+ __ocml_sin_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hcos(__half x)
+ {
+ return __half_raw{
+ __ocml_cos_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hexp(__half x)
+ {
+ return __half_raw{
+ __ocml_exp_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hexp2(__half x)
+ {
+ return __half_raw{
+ __ocml_exp2_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hexp10(__half x)
+ {
+ return __half_raw{
+ __ocml_exp10_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hlog2(__half x)
+ {
+ return __half_raw{
+ __ocml_log2_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hlog(__half x)
+ {
+ return __half_raw{
+ __ocml_log_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hlog10(__half x)
+ {
+ return __half_raw{
+ __ocml_log10_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hrcp(__half x)
+ {
+ return __half_raw{
+ static_cast<_Float16>(1.0f) /static_cast<__half_raw>(x).data};
+ }
+ inline
+ __device__
+ __half hrsqrt(__half x)
+ {
+ return __half_raw{
+ __ocml_rsqrt_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ __half hsqrt(__half x)
+ {
+ return __half_raw{
+ __ocml_sqrt_f16(static_cast<__half_raw>(x).data)};
+ }
+ inline
+ __device__
+ bool __hisinf(__half x)
+ {
+ return __ocml_isinf_f16(static_cast<__half_raw>(x).data);
+ }
+ inline
+ __device__
+ bool __hisnan(__half x)
+ {
+ return __ocml_isnan_f16(static_cast<__half_raw>(x).data);
+ }
+ inline
+ __device__
+ __half __hneg(__half x)
+ {
+ return __half_raw{-static_cast<__half_raw>(x).data};
+ }
+
+ inline
+ __HOST_DEVICE__
+ __half2 h2trunc(__half2 x)
+ {
+ return __half2{__ocml_trunc_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2ceil(__half2 x)
+ {
+ return __half2{__ocml_ceil_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2floor(__half2 x)
+ {
+ return __half2{__ocml_floor_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2rint(__half2 x)
+ {
+ return __half2{__ocml_rint_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2sin(__half2 x)
+ {
+ return __half2{__ocml_sin_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2cos(__half2 x)
+ {
+ return __half2{__ocml_cos_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2exp(__half2 x)
+ {
+ return __half2{__ocml_exp_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2exp2(__half2 x)
+ {
+ return __half2{__ocml_exp2_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2exp10(__half2 x)
+ {
+ return __half2{__ocml_exp10_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2log2(__half2 x)
+ {
+ return __half2{__ocml_log2_2f16(x)};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2log(__half2 x) { return __ocml_log_2f16(x); }
+ inline
+ __HOST_DEVICE__
+ __half2 h2log10(__half2 x) { return __ocml_log10_2f16(x); }
+ inline
+ __HOST_DEVICE__
+ __half2 h2rcp(__half2 x) {
+ return _Float16_2{
+ _Float16_2{static_cast<_Float16>(1.0f), static_cast<_Float16>(1.0f)} / x.data};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 h2rsqrt(__half2 x) { return __ocml_rsqrt_2f16(x); }
+ inline
+ __HOST_DEVICE__
+ __half2 h2sqrt(__half2 x) { return __ocml_sqrt_2f16(x); }
+ inline
+ __HOST_DEVICE__
+ __half2 __hisinf2(__half2 x)
+ {
+ auto r = __ocml_isinf_2f16(x);
+ return __half2{_Float16_2{
+ static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hisnan2(__half2 x)
+ {
+ auto r = __ocml_isnan_2f16(x);
+ return __half2{_Float16_2{
+ static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
+ }
+ inline
+ __HOST_DEVICE__
+ __half2 __hneg2(__half2 x)
+ {
+ return __half2{-static_cast<__half2_raw>(x).data};
+ }
+ } // Anonymous namespace.
+
+ #if !defined(HIP_NO_HALF)
+ using half = __half;
+ using half2 = __half2;
+ #endif
+ __device__
+ inline
+ __half __shfl(__half var, int src_lane, int width = warpSize) {
+ union { int i; __half h; } tmp; tmp.h = var;
+ tmp.i = __shfl(tmp.i, src_lane, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half2 __shfl(__half2 var, int src_lane, int width = warpSize) {
+ union { int i; __half2 h; } tmp; tmp.h = var;
+ tmp.i = __shfl(tmp.i, src_lane, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half __shfl_up(__half var, unsigned int lane_delta, int width = warpSize) {
+ union { int i; __half h; } tmp; tmp.h = var;
+ tmp.i = __shfl_up(tmp.i, lane_delta, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half2 __shfl_up(__half2 var, unsigned int lane_delta, int width = warpSize) {
+ union { int i; __half2 h; } tmp; tmp.h = var;
+ tmp.i = __shfl_up(tmp.i, lane_delta, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half __shfl_down(__half var, unsigned int lane_delta, int width = warpSize) {
+ union { int i; __half h; } tmp; tmp.h = var;
+ tmp.i = __shfl_down(tmp.i, lane_delta, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half2 __shfl_down(__half2 var, unsigned int lane_delta, int width = warpSize) {
+ union { int i; __half2 h; } tmp; tmp.h = var;
+ tmp.i = __shfl_down(tmp.i, lane_delta, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half __shfl_xor(__half var, int lane_mask, int width = warpSize) {
+ union { int i; __half h; } tmp; tmp.h = var;
+ tmp.i = __shfl_xor(tmp.i, lane_mask, width);
+ return tmp.h;
+ }
+ __device__
+ inline
+ __half2 __shfl_xor(__half2 var, int lane_mask, int width = warpSize) {
+ union { int i; __half2 h; } tmp; tmp.h = var;
+ tmp.i = __shfl_xor(tmp.i, lane_mask, width);
+ return tmp.h;
+ }
+ #endif // defined(__cplusplus)
+#elif defined(__GNUC__)
+ #if !defined(__HIPCC_RTC__)
+ #include "hip_fp16_gcc.h"
+ #endif
+#endif // !defined(__clang__) && defined(__GNUC__)
+
+#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP16_H
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_gl_interop.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_gl_interop.h
new file mode 100644
index 0000000000000000000000000000000000000000..e5b6dc3a359c6a5331721fbe3755d009eabd7a66
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_gl_interop.h
@@ -0,0 +1,108 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#ifndef HIP_INCLUDE_AMD_HIP_GL_INTEROP_H
+#define HIP_INCLUDE_AMD_HIP_GL_INTEROP_H
+
+#if defined(__cplusplus)
+extern "C" {
+#endif
+
+/**
+ *
+ * @addtogroup GlobalDefs
+ * @{
+ *
+ */
+
+/**
+ * HIP Devices used by current OpenGL Context.
+ */
+typedef enum hipGLDeviceList {
+ hipGLDeviceListAll = 1, ///< All hip devices used by current OpenGL context.
+ hipGLDeviceListCurrentFrame = 2, ///< Hip devices used by current OpenGL context in current
+ ///< frame
+ hipGLDeviceListNextFrame = 3 ///< Hip devices used by current OpenGL context in next
+ ///< frame.
+} hipGLDeviceList;
+
+
+/** GLuint as uint.*/
+typedef unsigned int GLuint;
+/** GLenum as uint.*/
+typedef unsigned int GLenum;
+/*
+* @}
+*/
+
+/**
+ * @ingroup GL
+ * @{
+ *
+ */
+/**
+ * @brief Queries devices associated with the current OpenGL context.
+ *
+ * @param [out] pHipDeviceCount - Pointer of number of devices on the current GL context.
+ * @param [out] pHipDevices - Pointer of devices on the current OpenGL context.
+ * @param [in] hipDeviceCount - Size of device.
+ * @param [in] deviceList - The setting of devices. It could be either hipGLDeviceListCurrentFrame
+ * for the devices used to render the current frame, or hipGLDeviceListAll for all devices.
+ * The default setting is Invalid deviceList value.
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
+ *
+ */
+hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices,
+ unsigned int hipDeviceCount, hipGLDeviceList deviceList);
+/**
+ * @brief Registers a GL Buffer for interop and returns corresponding graphics resource.
+ *
+ * @param [out] resource - Returns pointer of graphics resource.
+ * @param [in] buffer - Buffer to be registered.
+ * @param [in] flags - Register flags.
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorUnknown, #hipErrorInvalidResourceHandle
+ *
+ */
+hipError_t hipGraphicsGLRegisterBuffer(hipGraphicsResource** resource, GLuint buffer,
+ unsigned int flags);
+/**
+ * @brief Register a GL Image for interop and returns the corresponding graphic resource.
+ *
+ * @param [out] resource - Returns pointer of graphics resource.
+ * @param [in] image - Image to be registered.
+ * @param [in] target - Valid target value Id.
+ * @param [in] flags - Register flags.
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorUnknown, #hipErrorInvalidResourceHandle
+ *
+ */
+hipError_t hipGraphicsGLRegisterImage(hipGraphicsResource** resource, GLuint image,
+ GLenum target, unsigned int flags);
+/*
+* @}
+*/
+#if defined(__cplusplus)
+}
+#endif /* __cplusplus */
+#endif /* HIP_INCLUDE_AMD_HIP_GL_INTEROP_H */
diff --git a/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_math_constants.h b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_math_constants.h
new file mode 100644
index 0000000000000000000000000000000000000000..148fdcc4fb07f3b9fcca5bd6100bb6553413965a
--- /dev/null
+++ b/lib/python3.12/site-packages/triton/backends/amd/include/hip/amd_detail/amd_hip_math_constants.h
@@ -0,0 +1,124 @@
+/*
+Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#ifndef AMD_HIP_MATH_CONSTANTS_H
+#define AMD_HIP_MATH_CONSTANTS_H
+
+// single precision constants
+#define HIP_INF_F __int_as_float(0x7f800000U)
+#define HIP_NAN_F __int_as_float(0x7fffffffU)
+#define HIP_MIN_DENORM_F __int_as_float(0x00000001U)
+#define HIP_MAX_NORMAL_F __int_as_float(0x7f7fffffU)
+#define HIP_NEG_ZERO_F __int_as_float(0x80000000U)
+#define HIP_ZERO_F 0.0F
+#define HIP_ONE_F 1.0F
+#define HIP_SQRT_HALF_F 0.707106781F
+#define HIP_SQRT_HALF_HI_F 0.707106781F
+#define HIP_SQRT_HALF_LO_F 1.210161749e-08F
+#define HIP_SQRT_TWO_F 1.414213562F
+#define HIP_THIRD_F 0.333333333F
+#define HIP_PIO4_F 0.785398163F
+#define HIP_PIO2_F 1.570796327F
+#define HIP_3PIO4_F 2.356194490F
+#define HIP_2_OVER_PI_F 0.636619772F
+#define HIP_SQRT_2_OVER_PI_F 0.797884561F
+#define HIP_PI_F 3.141592654F
+#define HIP_L2E_F 1.442695041F
+#define HIP_L2T_F 3.321928094F
+#define HIP_LG2_F 0.301029996F
+#define HIP_LGE_F 0.434294482F
+#define HIP_LN2_F 0.693147181F
+#define HIP_LNT_F 2.302585093F
+#define HIP_LNPI_F 1.144729886F
+#define HIP_TWO_TO_M126_F 1.175494351e-38F
+#define HIP_TWO_TO_126_F 8.507059173e37F
+#define HIP_NORM_HUGE_F 3.402823466e38F
+#define HIP_TWO_TO_23_F 8388608.0F
+#define HIP_TWO_TO_24_F 16777216.0F
+#define HIP_TWO_TO_31_F 2147483648.0F
+#define HIP_TWO_TO_32_F 4294967296.0F
+#define HIP_REMQUO_BITS_F 3U
+#define HIP_REMQUO_MASK_F (~((~0U)<
+
+#if !defined(__HIPCC_RTC__)
+#ifdef __cplusplus
+#include
+#else
+#include
+#endif // __cplusplus
+#endif // !defined(__HIPCC_RTC__)
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/**
+ * @brief Query the installed library build name.
+ *
+ * This function can be used even when the library is not initialized.
+ *
+ * @returns Returns a string describing the build version of the library. The
+ * string is owned by the library.
+ */
+const char* amd_dbgapi_get_build_name();
+
+/**
+ * @brief Query the installed library git hash.
+ *
+ * This function can be used even when the library is not initialized.
+ *
+ * @returns Returns git hash of the library.
+ */
+const char* amd_dbgapi_get_git_hash();
+
+/**
+ * @brief Query the installed library build ID.
+ *
+ * This function can be used even when the library is not initialized.
+ *
+ * @returns Returns build ID of the library.
+ */
+size_t amd_dbgapi_get_build_id();
+
+#ifdef __cplusplus
+} /* extern "c" */
+#endif
+
+//---
+// Top part of file can be compiled with any compiler
+
+#if !defined(__HIPCC_RTC__)
+#ifdef __cplusplus
+#include
+#include
+#include
+#else
+#include
+#include
+#endif // __cplusplus
+#else
+#if !__HIP_NO_STD_DEFS__
+typedef unsigned int uint32_t;
+typedef unsigned long long uint64_t;
+typedef signed int int32_t;
+typedef signed long long int64_t;
+namespace std {
+using ::uint32_t;
+using ::uint64_t;
+using ::int32_t;
+using ::int64_t;
+}
+#endif // __HIP_NO_STD_DEFS__
+#endif // !defined(__HIPCC_RTC__)
+
+#if __HIP_CLANG_ONLY__
+
+#if !defined(__align__)
+#define __align__(x) __attribute__((aligned(x)))
+#endif
+
+#define CUDA_SUCCESS hipSuccess
+
+#if !defined(__HIPCC_RTC__)
+#include
+#include
+#include
+#include
+#include
+#include
+extern int HIP_TRACE_API;
+#endif // !defined(__HIPCC_RTC__)
+
+#ifdef __cplusplus
+#include
+#endif
+
+#include
+
+// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
+#if defined(__KALMAR_ACCELERATOR__) && !defined(__HCC_ACCELERATOR__)
+#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
+#endif
+
+// Feature tests:
+#if (defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)) || __HIP_DEVICE_COMPILE__
+// Device compile and not host compile:
+
+// 32-bit Atomics:
+#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (1)
+#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (1)
+#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (1)
+#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (1)
+#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (1)
+
+// 64-bit Atomics:
+#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (1)
+#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (1)
+
+// Doubles
+#define __HIP_ARCH_HAS_DOUBLES__ (1)
+
+// warp cross-lane operations:
+#define __HIP_ARCH_HAS_WARP_VOTE__ (1)
+#define __HIP_ARCH_HAS_WARP_BALLOT__ (1)
+#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (1)
+#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
+
+// sync
+#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (1)
+#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
+
+// misc
+#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
+#define __HIP_ARCH_HAS_3DGRID__ (1)
+#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
+
+#endif /* Device feature flags */
+
+
+#define launch_bounds_impl0(requiredMaxThreadsPerBlock) \
+ __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock)))
+#define launch_bounds_impl1(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor) \
+ __attribute__((amdgpu_flat_work_group_size(1, requiredMaxThreadsPerBlock), \
+ amdgpu_waves_per_eu(minBlocksPerMultiprocessor)))
+#define select_impl_(_1, _2, impl_, ...) impl_
+#define __launch_bounds__(...) \
+ select_impl_(__VA_ARGS__, launch_bounds_impl1, launch_bounds_impl0, )(__VA_ARGS__)
+
+#if !defined(__HIPCC_RTC__)
+__host__ inline void* __get_dynamicgroupbaseptr() { return nullptr; }
+#endif // !defined(__HIPCC_RTC__)
+
+// End doxygen API:
+/**
+ * @}
+ */
+
+//
+// hip-clang functions
+//
+#if !defined(__HIPCC_RTC__)
+#define HIP_KERNEL_NAME(...) __VA_ARGS__
+#define HIP_SYMBOL(X) X
+
+typedef int hipLaunchParm;
+
+template ::type* = nullptr>
+void pArgs(const std::tuple&, void*) {}
+
+template ::type* = nullptr>
+void pArgs(const std::tuple& formals, void** _vargs) {
+ using T = typename std::tuple_element >::type;
+
+ static_assert(!std::is_reference{},
+ "A __global__ function cannot have a reference as one of its "
+ "arguments.");
+#if defined(HIP_STRICT)
+ static_assert(std::is_trivially_copyable{},
+ "Only TriviallyCopyable types can be arguments to a __global__ "
+ "function");
+#endif
+ _vargs[n] = const_cast(reinterpret_cast(&std::get