diff options
Diffstat (limited to 'llvm/docs')
| -rw-r--r-- | llvm/docs/NVPTXUsage.rst | 444 | ||||
| -rw-r--r-- | llvm/docs/Reference.rst | 4 | ||||
| -rw-r--r-- | llvm/docs/Security.rst | 2 | ||||
| -rw-r--r-- | llvm/docs/YamlIO.rst | 92 |
4 files changed, 495 insertions, 47 deletions
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst index e8dceb8..5ad8f9a 100644 --- a/llvm/docs/NVPTXUsage.rst +++ b/llvm/docs/NVPTXUsage.rst @@ -322,6 +322,450 @@ aligned '``@llvm.nvvm.barrier.cta.*``' instruction should only be used if it is known that all threads in the CTA evaluate the condition identically, otherwise behavior is undefined. +MBarrier family of Intrinsics +----------------------------- + +Overview: +^^^^^^^^^ + +An ``mbarrier`` is a barrier created in shared memory that supports: + +* Synchronizing any subset of threads within a CTA. +* One-way synchronization of threads across CTAs of a cluster. + Threads can perform only ``arrive`` operations but not ``*_wait`` on an + mbarrier located in shared::cluster space. +* Waiting for completion of asynchronous memory operations initiated by a + thread and making them visible to other threads. + +Unlike ``bar{.cta}/barrier{.cta}`` instructions which can access a limited +number of barriers per CTA, ``mbarrier`` objects are user-defined and are +only limited by the total shared memory size available. + +An mbarrier object is an opaque object in shared memory with an +alignment of 8-bytes. It keeps track of: + +* Current phase of the mbarrier object +* Count of pending arrivals for the current phase of the mbarrier object +* Count of expected arrivals for the next phase of the mbarrier object +* Count of pending asynchronous memory operations (or transactions) + tracked by the current phase of the mbarrier object. This is also + referred to as ``tx-count``. The unit of ``tx-count`` is specified + by the asynchronous memory operation (for example, + ``llvm.nvvm.cp.async.bulk.tensor.g2s.*``). + +The ``phase`` of an mbarrier object is the number of times the mbarrier +object has been used to synchronize threads/track async operations. +In each phase, threads perform: + +* arrive/expect-tx/complete-tx operations to progress the current phase. +* test_wait/try_wait operations to check for completion of the current phase. + +An mbarrier object completes the current phase when: + +* The count of the pending arrivals has reached zero AND +* The tx-count has reached zero. + +When an mbarrier object completes the current phase, below +actions are performed ``atomically``: + +* The mbarrier object transitions to the next phase. +* The pending arrival count is reinitialized to the expected arrival count. + +For more information, refer PTX ISA +`<https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier>`_. + +'``llvm.nvvm.mbarrier.init``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare void @llvm.nvvm.mbarrier.init(ptr %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3) %addr, i32 %count) + +Overview: +""""""""" + +The '``@llvm.nvvm.mbarrier.init.*``' intrinsics are used to initialize +an mbarrier object located at ``addr`` with the value ``count``. +``count`` is a 32-bit unsigned integer value and must be within +the range [1...2^20-1]. During initialization: + +* The tx-count and the current phase of the mbarrier object are set to 0. +* The expected and pending arrival counts are set to ``count``. + +Semantics: +"""""""""" + +The ``.shared`` variant explicitly uses shared memory address space for +the ``addr`` operand. If the ``addr`` does not fall within the +shared::cta space, then the behavior of this intrinsic is undefined. +Performing ``mbarrier.init`` on a valid mbarrier object is undefined; +use ``mbarrier.inval`` before reusing the memory for another mbarrier +or any other purpose. + +'``llvm.nvvm.mbarrier.inval``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare void @llvm.nvvm.mbarrier.inval(ptr %addr) + declare void @llvm.nvvm.mbarrier.inval.shared(ptr addrspace(3) %addr) + +Overview: +""""""""" + +The '``@llvm.nvvm.mbarrier.inval.*``' intrinsics invalidate the mbarrier +object at the address specified by ``addr``. + +Semantics: +"""""""""" + +The ``.shared`` variant explicitly uses shared memory address space for +the ``addr`` operand. If the ``addr`` does not fall within the +shared::cta space, then the behavior of this intrinsic is undefined. +It is expected that ``addr`` was previously initialized using +``mbarrier.init``; otherwise, the behavior is undefined. + +'``llvm.nvvm.mbarrier.expect.tx``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + +Overview: +""""""""" + +The '``@llvm.nvvm.mbarrier.expect.tx.*``' intrinsics increase the transaction +count of the mbarrier object at ``%addr`` by ``%tx_count``. The ``%tx_count`` +is a 32-bit unsigned integer value. + +Semantics: +"""""""""" + +The ``.space.{cta/cluster}`` indicates the address space where the mbarrier +object resides. + +The ``.scope.{cta/cluster}`` denotes the set of threads that can directly +observe the synchronizing effect of the mbarrier operation. When scope is +"cta", all threads executing in the same CTA (as the current thread) can +directly observe the effect of the ``expect.tx`` operation. Similarly, +when scope is "cluster", all threads executing in the same Cluster +(as the current thread) can directly observe the effect of the operation. + +If the ``addr`` does not fall within shared::cta or shared::cluster space, +then the behavior of this intrinsic is undefined. This intrinsic has +``relaxed`` semantics and hence does not provide any memory ordering +or visibility guarantees. + +'``llvm.nvvm.mbarrier.complete.tx``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.complete.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.complete.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + +Overview: +""""""""" + +The '``@llvm.nvvm.mbarrier.complete.tx.*``' intrinsics decrease the transaction +count of the mbarrier object at ``%addr`` by ``%tx_count``. The ``%tx_count`` +is a 32-bit unsigned integer value. As a result of this decrement, +the mbarrier can potentially complete its current phase and transition +to the next phase. + +Semantics: +"""""""""" + +The semantics of these intrinsics are identical to those of the +``llvm.nvvm.mbarrier.expect.tx.*`` intrinsics described above. + +'``llvm.nvvm.mbarrier.arrive``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare i64 @llvm.nvvm.mbarrier.arrive.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count) + declare i64 @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count) + + declare i64 @llvm.nvvm.mbarrier.arrive.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count) + declare i64 @llvm.nvvm.mbarrier.arrive.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count) + +Overview: +""""""""" + +The ``@llvm.nvvm.mbarrier.arrive.*`` intrinsics signal the arrival of the +executing thread or completion of an asynchronous instruction associated with +an arrive operation on the mbarrier object at ``%addr``. This operation +decrements the pending arrival count by ``%count``, a 32-bit unsigned integer, +potentially completing the current phase and triggering a transition to the +next phase. + +Semantics: +"""""""""" + +The ``.space.{cta/cluster}`` indicates the address space where the mbarrier +object resides. When the mbarrier is in shared::cta space, the intrinsics +return an opaque 64-bit value capturing the phase of the mbarrier object +_prior_ to this arrive operation. This value can be used with a try_wait +or test_wait operation to check for the completion of the mbarrier. + +The ``.scope.{cta/cluster}`` denotes the set of threads that can directly +observe the synchronizing effect of the mbarrier operation. When scope is +"cta", all threads executing in the same CTA (as the current thread) can +directly observe the effect of the ``arrive`` operation. Similarly, +when scope is "cluster", all threads executing in the same Cluster +(as the current thread) can directly observe the effect of the operation. + +If the ``addr`` does not fall within shared::cta or shared::cluster space, +then the behavior of this intrinsic is undefined. + +These intrinsics have ``release`` semantics by default. The release semantics +ensure ordering of operations that occur in program order _before_ this arrive +instruction, making their effects visible to subsequent operations in other +threads of the CTA (or cluster, depending on scope). Threads performing +corresponding acquire operations (such as mbarrier.test.wait) synchronize +with this release. The ``relaxed`` variants of these intrinsics do not +provide any memory ordering or visibility guarantees. + +'``llvm.nvvm.mbarrier.arrive.expect.tx``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare i64 @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare i64 @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + + declare i64 @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare i64 @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.expect.tx.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + +Overview: +""""""""" + +The ``@llvm.nvvm.mbarrier.arrive.expect.tx.*`` intrinsics are similar to +the ``@llvm.nvvm.mbarrier.arrive`` intrinsics except that they also +perform an ``expect-tx`` operation _prior_ to the ``arrive`` operation. +The ``%tx_count`` specifies the transaction count for the ``expect-tx`` +operation and the count for the ``arrive`` operation is assumed to be 1. + +Semantics: +"""""""""" + +The semantics of these intrinsics are identical to those of the +``llvm.nvvm.mbarrier.arrive.*`` intrinsics described above. + +'``llvm.nvvm.mbarrier.arrive.drop``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare i64 @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count) + declare i64 @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.drop.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.drop.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count) + + declare i64 @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %count) + declare i64 @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %count) + declare void @llvm.nvvm.mbarrier.arrive.drop.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %count) + +Overview: +""""""""" + +The ``@llvm.nvvm.mbarrier.arrive.drop.*`` intrinsics decrement the +expected arrival count of the mbarrier object at ``%addr`` by +``%count`` and then perform an ``arrive`` operation with ``%count``. +The ``%count`` is a 32-bit integer. + +Semantics: +"""""""""" + +The semantics of these intrinsics are identical to those of the +``llvm.nvvm.mbarrier.arrive.*`` intrinsics described above. + +'``llvm.nvvm.mbarrier.arrive.drop.expect.tx``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + + declare i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare i64 @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cta.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + declare void @llvm.nvvm.mbarrier.arrive.drop.expect.tx.relaxed.scope.cluster.space.cluster(ptr addrspace(7) %addr, i32 %tx_count) + +Overview: +""""""""" + +The ``@llvm.nvvm.mbarrier.arrive.drop.expect.tx.*`` intrinsics perform +the below operations on the mbarrier located at ``%addr``. + +* Perform an ``expect-tx`` operation i.e. increase the transaction count + of the mbarrier by ``%tx_count``, a 32-bit unsigned integer value. +* Decrement the expected arrival count of the mbarrier by 1. +* Perform an ``arrive`` operation on the mbarrier with a value of 1. + +Semantics: +"""""""""" + +The semantics of these intrinsics are identical to those of the +``llvm.nvvm.mbarrier.arrive.*`` intrinsics described above. + +'``llvm.nvvm.mbarrier.test.wait``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare i1 @llvm.nvvm.mbarrier.test.wait.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state) + declare i1 @llvm.nvvm.mbarrier.test.wait.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state) + declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase) + declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase) + + declare i1 @llvm.nvvm.mbarrier.test.wait.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state) + declare i1 @llvm.nvvm.mbarrier.test.wait.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state) + declare i1 @llvm.nvvm.mbarrier.test.wait.parity.relaxed.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase) + declare i1 @llvm.nvvm.mbarrier.test.wait.parity.relaxed.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase) + +Overview: +""""""""" + +The ``@llvm.nvvm.mbarrier.test.wait.*`` intrinsics test for the completion +of the current or the immediately preceding phase of an mbarrier object at +``%addr``. The test for completion can be done with either the ``state`` or +the ``phase-parity`` of the mbarrier object. + +* When done through the ``i64 %state`` operand, the state must be + returned by an ``llvm.nvvm.mbarrier.arrive.*`` on the _same_ + mbarrier object. +* The ``.parity`` variant of these intrinsics test for completion + of the phase indicated by the operand ``i32 %phase``, which is + the integer parity of either the current phase or the immediately + preceding phase of the mbarrier object. An even phase has integer + parity 0 and an odd phase has integer parity of 1. So the valid + values for phase-parity are 0 and 1. + +Semantics: +"""""""""" + +The ``.scope.{cta/cluster}`` denotes the set of threads that the +test_wait operation can directly synchronize with. + +If the ``addr`` does not fall within shared::cta space, then the +the behavior of this intrinsic is undefined. + +These intrinsics have ``acquire`` semantics by default. This acquire +pattern establishes memory ordering for operations occurring in program +order after this ``test_wait`` instruction by making operations from +other threads in the CTA (or cluster, depending on scope) visible to +subsequent operations in the current thread. When this wait completes, +it synchronizes with the corresponding release pattern from the +``mbarrier.arrive`` operation. The ``relaxed`` variants of these intrinsics +do not provide any memory ordering or visibility guarantees. + +This ``test.wait`` intrinsic is non-blocking and immediately returns +the completion status without suspending the executing thread. + +The boolean return value indicates: + +* True: The immediately preceding phase has completed +* False: The current phase is still incomplete + +When this wait returns true, the following ordering guarantees hold: + +* All memory accesses (except async operations) requested prior to + ``mbarrier.arrive`` having release semantics by participating + threads of a CTA (or cluster, depending on scope) are visible to + the executing thread. +* All ``cp.async`` operations requested prior to ``cp.async.mbarrier.arrive`` + by participating threads of a CTA are visible to the executing thread. +* All ``cp.async.bulk`` operations using the same mbarrier object requested + prior to ``mbarrier.arrive`` having release semantics by participating CTA + threads are visible to the executing thread. +* Memory accesses requested after this wait are not visible to memory + accesses performed prior to ``mbarrier.arrive`` by other participating + threads. +* No ordering guarantee exists for memory accesses by the same thread + between an ``mbarrier.arrive`` and this wait. + +'``llvm.nvvm.mbarrier.try.wait``' +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Syntax: +""""""" + +.. code-block:: llvm + + declare i1 @llvm.nvvm.mbarrier.try.wait{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state) + declare i1 @llvm.nvvm.mbarrier.try.wait{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state) + + declare i1 @llvm.nvvm.mbarrier.try.wait.parity{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase) + declare i1 @llvm.nvvm.mbarrier.try.wait.parity{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase) + + declare i1 @llvm.nvvm.mbarrier.try.wait.tl{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i64 %state, i32 %timelimit) + declare i1 @llvm.nvvm.mbarrier.try.wait.tl{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i64 %state, i32 %timelimit) + + declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl{.relaxed}.scope.cta.space.cta(ptr addrspace(3) %addr, i32 %phase, i32 %timelimit) + declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl{.relaxed}.scope.cluster.space.cta(ptr addrspace(3) %addr, i32 %phase, i32 %timelimit) + +Overview: +""""""""" + +The ``@llvm.nvvm.mbarrier.try.wait.*`` intrinsics test for the completion of +the current or immediately preceding phase of an mbarrier object at ``%addr``. +Unlike the ``test.wait`` intrinsics, which perform a non-blocking test, these +intrinsics may block the executing thread until the specified phase completes +or a system-dependent time limit expires. Suspended threads resume execution +when the phase completes or the time limit elapses. This time limit is +configurable through the ``.tl`` variants of these intrinsics, where the +``%timelimit`` operand (an unsigned integer) specifies the limit in +nanoseconds. Other semantics are identical to those of the ``test.wait`` +intrinsics described above. + Electing a thread ----------------- diff --git a/llvm/docs/Reference.rst b/llvm/docs/Reference.rst index 9b1bf1b..212aefa 100644 --- a/llvm/docs/Reference.rst +++ b/llvm/docs/Reference.rst @@ -250,3 +250,7 @@ Additional Topics :doc:`ContentAddressableStorage` A reference guide for using LLVM's CAS library. + +:doc:`CIBestPractices` + A list of guidelines and best practices to use when working on LLVM's + CI systems. diff --git a/llvm/docs/Security.rst b/llvm/docs/Security.rst index 5cb8d04..d95b552 100644 --- a/llvm/docs/Security.rst +++ b/llvm/docs/Security.rst @@ -51,7 +51,7 @@ username for an individual isn't available, the brackets will be empty. * Nikhil Gupta (Nvidia) [] * Oliver Hunt (Apple) [@ojhunt] * Peter Smith (ARM) [@smithp35] -* Pietro Albini (Ferrous Systems; Rust) [@pietroalbini] +* Pietro Albini (Oxide Computer Company; Rust) [@pietroalbini] * Serge Guelton (Mozilla) [@serge-sans-paille] * Sergey Zverev (Intel) [@offsake] * Shayne Hiet-Block (Microsoft) [@GreatKeeper] diff --git a/llvm/docs/YamlIO.rst b/llvm/docs/YamlIO.rst index c5079d8..553e4a8 100644 --- a/llvm/docs/YamlIO.rst +++ b/llvm/docs/YamlIO.rst @@ -12,7 +12,7 @@ YAML is a human-readable data serialization language. The full YAML language spec can be read at `yaml.org <http://www.yaml.org/spec/1.2/spec.html#Introduction>`_. The simplest form of YAML is just "scalars", "mappings", and "sequences". A scalar is any number -or string. The pound/hash symbol (#) begins a comment line. A mapping is +or string. The pound/hash symbol (``#``) begins a comment line. A mapping is a set of key-value pairs where the key ends with a colon. For example: .. code-block:: yaml @@ -21,7 +21,7 @@ a set of key-value pairs where the key ends with a colon. For example: name: Tom hat-size: 7 -A sequence is a list of items where each item starts with a leading dash ('-'). +A sequence is a list of items where each item starts with a leading dash (``-``). For example: .. code-block:: yaml @@ -221,7 +221,7 @@ Similar errors are produced for other input not conforming to the schema. Scalars ======= -YAML scalars are just strings (i.e. not a sequence or mapping). The YAML I/O +YAML scalars are just strings (i.e., not a sequence or mapping). The YAML I/O library provides support for translating between YAML scalars and specific C++ types. @@ -230,19 +230,19 @@ Built-in types -------------- The following types have built-in support in YAML I/O: -* bool -* float -* double -* StringRef -* std::string -* int64_t -* int32_t -* int16_t -* int8_t -* uint64_t -* uint32_t -* uint16_t -* uint8_t +* ``bool`` +* ``float`` +* ``double`` +* ``StringRef`` +* ``std::string`` +* ``int64_t`` +* ``int32_t`` +* ``int16_t`` +* ``int8_t`` +* ``uint64_t`` +* ``uint32_t`` +* ``uint16_t`` +* ``uint8_t`` That is, you can use those types in fields of ``MappingTraits`` or as the element type in a sequence. When reading, YAML I/O will validate that the string found @@ -264,7 +264,7 @@ operators to and from the base type. For example: LLVM_YAML_STRONG_TYPEDEF(uint32_t, MyFooFlags) LLVM_YAML_STRONG_TYPEDEF(uint32_t, MyBarFlags) -This generates two classes MyFooFlags and MyBarFlags which you can use in your +This generates two classes ``MyFooFlags`` and ``MyBarFlags`` which you can use in your native data structures instead of ``uint32_t``. They are implicitly converted to and from ``uint32_t``. The point of creating these unique types is that you can now specify traits on them to get different YAML conversions. @@ -275,12 +275,12 @@ An example use of a unique type is that YAML I/O provides fixed-sized unsigned integers that are written with YAML I/O as hexadecimal instead of the decimal format used by the built-in integer types: -* Hex64 -* Hex32 -* Hex16 -* Hex8 +* ``Hex64`` +* ``Hex32`` +* ``Hex16`` +* ``Hex8`` -You can use ``llvm::yaml::Hex32`` instead of ``uint32_t`` and the only difference will +You can use ``llvm::yaml::Hex32`` instead of ``uint32_t``. The only difference will be that when YAML I/O writes out that type it will be formatted in hexadecimal. @@ -356,8 +356,8 @@ had the following bit flags defined: LLVM_YAML_STRONG_TYPEDEF(uint32_t, MyFlags) -To support reading and writing of MyFlags, you specialize ``ScalarBitSetTraits<>`` -on MyFlags and provide the bit values and their names. +To support reading and writing of ``MyFlags``, you specialize ``ScalarBitSetTraits<>`` +on ``MyFlags`` and provide the bit values and their names. .. code-block:: c++ @@ -441,7 +441,7 @@ Custom Scalar Sometimes, for readability, a scalar needs to be formatted in a custom way. For instance, your internal data structure may use an integer for time (seconds since some epoch), but in YAML it would be much nicer to express that integer in -some time format (e.g. 4-May-2012 10:30pm). YAML I/O has a way to support +some time format (e.g., ``4-May-2012 10:30pm``). YAML I/O has a way to support custom formatting and parsing of scalar types by specializing ``ScalarTraits<>`` on your data type. When writing, YAML I/O will provide the native type and your specialization must create a temporary ``llvm::StringRef``. When reading, @@ -523,8 +523,8 @@ An example of a custom type with an appropriate specialization of Mappings ======== -To be translated to or from a YAML mapping for your type T you must specialize -``llvm::yaml::MappingTraits`` on T and implement the "void mapping(IO &io, T&)" +To be translated to or from a YAML mapping for your type ``T``, you must specialize +``llvm::yaml::MappingTraits`` on ``T`` and implement the ``void mapping(IO &io, T&)`` method. If your native data structures use pointers to a class everywhere, you can specialize on the class pointer. Examples: @@ -685,13 +685,13 @@ normalized instance is stack allocated. In these cases, the utility template ``MappingNormalizationHeap<>`` can be used instead. It just like ``MappingNormalization<>`` except that it heap allocates the normalized object when reading YAML. It never destroys the normalized object. The ``denormalize()`` -method can this return "this". +method can this return ``this``. Default values -------------- Within a ``mapping()`` method, calls to ``io.mapRequired()`` mean that that key is -required to exist when parsing YAML documents, otherwise YAML I/O will issue an +required to exist when parsing YAML documents; otherwise, YAML I/O will issue an error. On the other hand, keys registered with ``io.mapOptional()`` are allowed to not @@ -708,7 +708,7 @@ does not have that key. There is one important difference between those two ways (default constructor and third parameter to ``mapOptional()``). When YAML I/O generates a YAML document, if the ``mapOptional()`` third parameter is used, if the actual value being written -is the same as (using ==) the default value, then that key/value is not written. +is the same as (using ``==``) the default value, then that key/value is not written. Order of Keys @@ -772,7 +772,7 @@ not. This is similar to something having no syntax errors, but still having semantic errors. To support semantic-level checking, YAML I/O allows an optional ``validate()`` method in a MappingTraits template specialization. -When parsing YAML, the ``validate()`` method is call *after* all key/values in +When parsing YAML, the ``validate()`` method is called *after* all key/values in the map have been processed. Any error message returned by the ``validate()`` method during input will be printed just like a syntax error would be printed. When writing YAML, the ``validate()`` method is called *before* the YAML @@ -807,7 +807,7 @@ Flow Mapping A YAML "flow mapping" is a mapping that uses the inline notation (e.g { x: 1, y: 0 } ) when written to YAML. To specify that a type should be written in YAML using flow mapping, your MappingTraits specialization should -add "static const bool flow = true;". For instance: +add ``static constexpr bool flow = true;``. For instance: .. code-block:: c++ @@ -824,7 +824,7 @@ add "static const bool flow = true;". For instance: ... } - static const bool flow = true; + static constexpr bool flow = true; } Flow mappings are subject to line wrapping according to the ``Output`` object @@ -833,8 +833,8 @@ configuration. Sequence ======== -To be translated to or from a YAML sequence for your type T you must specialize -``llvm::yaml::SequenceTraits`` on T and implement two methods: +To be translated to or from a YAML sequence for your type ``T``, you must specialize +``llvm::yaml::SequenceTraits`` on ``T`` and implement two methods: ``size_t size(IO &io, T&)`` and ``T::value_type& element(IO &io, T&, size_t indx)``. For example: @@ -857,9 +857,9 @@ a reference to that new allocated space. Flow Sequence ------------- A YAML "flow sequence" is a sequence that when written to YAML it uses the -inline notation (e.g [ foo, bar ] ). To specify that a sequence type should +inline notation (e.g., ``[ foo, bar ]`` ). To specify that a sequence type should be written in YAML as a flow sequence, your SequenceTraits specialization should -add "static const bool flow = true;". For instance: +add ``static constexpr bool flow = true;``. For instance: .. code-block:: c++ @@ -869,12 +869,12 @@ add "static const bool flow = true;". For instance: static MyListEl &element(IO &io, MyList &list, size_t index) { ... } // The existence of this member causes YAML I/O to use a flow sequence - static const bool flow = true; + static constexpr bool flow = true; }; -With the above, if you used MyList as the data type in your native data +With the above, if you used ``MyList`` as the data type in your native data structures, then when converted to YAML, a flow sequence of integers -will be used (e.g. [ 10, -3, 4 ]). +will be used (e.g., ``[ 10, -3, 4 ]``). Flow sequences are subject to line wrapping according to the Output object configuration. @@ -900,8 +900,8 @@ Document List ============= YAML allows you to define multiple "documents" in a single YAML file. Each -new document starts with a left aligned "---" token. The end of all documents -is denoted with a left aligned "..." token. Many users of YAML will never +new document starts with a left aligned ``---`` token. The end of all documents +is denoted with a left aligned ``...`` token. Many users of YAML will never have need for multiple documents. The top level node in their YAML schema will be a mapping or sequence. For those cases, the following is not needed. But for cases where you do want multiple documents, you can specify a @@ -955,7 +955,7 @@ to write your native data as YAML. One thing to recall is that a YAML file can contain multiple "documents". If the top level data structure you are streaming as YAML is a mapping, scalar, or sequence, then ``Output`` assumes you are generating one document and wraps the mapping output -with "``---``" and trailing "``...``". +with ``---`` and trailing ``...``. The ``WrapColumn`` parameter will cause the flow mappings and sequences to line-wrap when they go over the supplied column. Pass 0 to completely @@ -981,8 +981,8 @@ The above could produce output like: On the other hand, if the top level data structure you are streaming as YAML has a ``DocumentListTraits`` specialization, then Output walks through each element -of your DocumentList and generates a "---" before the start of each element -and ends with a "...". +of your DocumentList and generates a ``---`` before the start of each element +and ends with a ``...``. .. code-block:: c++ @@ -1022,7 +1022,7 @@ pointer: Once you have an ``Input`` object, you can use the C++ stream operator to read the document(s). If you expect there might be multiple YAML documents in one file, you'll need to specialize ``DocumentListTraits`` on a list of your -document type and stream in that document list type. Otherwise you can +document type and stream in that document list type. Otherwise, you can just stream in the document type. Also, you can check if there was any syntax errors in the YAML by calling the ``error()`` method on the ``Input`` object. For example: |