diff --git a/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst b/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst index 41528ec96b4b..405ccf0e0a9e 100644 --- a/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst +++ b/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait.rst @@ -134,3 +134,75 @@ mbarrier.test_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 bool& isReportSeen, uint64_t* addr, uint64_t state); + +mbarrier.test_wait.phase_type::primary.acquire.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + +mbarrier.test_wait.phase_type::primary.acquire.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + +mbarrier.test_wait.phase_type::primary.relaxed.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + +mbarrier.test_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); diff --git a/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst b/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst index be087650f1ef..0dd02b907695 100644 --- a/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst +++ b/docs/libcudacxx/ptx/instructions/generated/mbarrier_test_wait_parity.rst @@ -135,6 +135,79 @@ mbarrier.test_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 uint64_t* addr, uint32_t phaseParity); + +mbarrier.test_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + +mbarrier.test_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + +mbarrier.test_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + +mbarrier.test_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_test_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + mbarrier.test_wait.parity.phase_type::conditional.acquire.cta.shared::cta.b64 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: cuda diff --git a/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst b/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst index b436057bcab6..c35f1cfcf9e6 100644 --- a/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst +++ b/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait.rst @@ -206,6 +206,79 @@ mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 uint64_t* addr, uint64_t state); + +mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + +mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + +mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + +mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); + mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: cuda @@ -277,3 +350,79 @@ mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 uint64_t* addr, uint64_t state, uint32_t suspendTimeHint); + +mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state, + uint32_t suspendTimeHint); + +mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state, + uint32_t suspendTimeHint); + +mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state, + uint32_t suspendTimeHint); + +mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state, + uint32_t suspendTimeHint); diff --git a/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst b/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst index b339deaf1aea..80cc8c75759e 100644 --- a/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst +++ b/docs/libcudacxx/ptx/instructions/generated/mbarrier_try_wait_parity.rst @@ -206,6 +206,79 @@ mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 uint64_t* addr, uint32_t phaseParity); + +mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + +mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + +mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + +mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); + mbarrier.try_wait.parity.phase_type::conditional.acquire.cta.shared::cta.b64 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: cuda @@ -342,6 +415,83 @@ mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 uint32_t phaseParity, uint32_t suspendTimeHint); + +mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity, + uint32_t suspendTimeHint); + +mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity, + uint32_t suspendTimeHint); + +mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity, + uint32_t suspendTimeHint); + +mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +.. code-block:: cuda + + // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90 + // .phase_type = { .phase_type::primary } + // .sem = { .acquire, .relaxed } + // .scope = { .cta, .cluster } + template + __device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity, + uint32_t suspendTimeHint); + mbarrier.try_wait.parity.phase_type::conditional.acquire.cta.shared::cta.b64 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: cuda diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h index b2c9d36d06a6..7ea40bcfc0a0 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait.h @@ -217,6 +217,113 @@ _CCCL_DEVICE static inline bool mbarrier_test_wait( } #endif // __cccl_ptx_isa >= 940 +/* +// mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX +ISA 94, SM_90 +// .phase_type = { .phase_type::primary } +// .sem = { .acquire, .relaxed } +// .scope = { .cta, .cluster } +template +__device__ static inline bool mbarrier_test_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); +*/ +#if __cccl_ptx_isa >= 940 +template <::cuda::ptx::dot_sem _Sem, ::cuda::ptx::dot_scope _Scope> +_CCCL_DEVICE static inline bool mbarrier_test_wait( + ::cuda::ptx::mbarrier_phase_primary_t, + ::cuda::ptx::sem_t<_Sem> __sem, + ::cuda::ptx::scope_t<_Scope> __scope, + bool& __isReportSeen, + ::cuda::std::uint8_t& __reportValue, + ::cuda::std::uint64_t* __addr, + ::cuda::std::uint64_t __state) +{ + // __phase_type == mbarrier_phase_primary (due to parameter type constraint) + static_assert(__sem == sem_acquire || __sem == sem_relaxed, ""); + static_assert(__scope == scope_cta || __scope == scope_cluster, ""); + ::cuda::std::uint32_t __waitComplete; + ::cuda::std::uint32_t __isReportSeen_tmp; + ::cuda::std::uint32_t __reportValue_tmp; + if constexpr (__sem == sem_acquire && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.phase_type::primary.acquire.cta.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + else if constexpr (__sem == sem_acquire && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.phase_type::primary.acquire.cluster.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.phase_type::primary.relaxed.cta.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + __isReportSeen = static_cast(__isReportSeen_tmp); + __reportValue = static_cast<::cuda::std::uint8_t>(__reportValue_tmp); + return static_cast(__waitComplete); +} +#endif // __cccl_ptx_isa >= 940 + // NOLINTEND(modernize-unary-static-assert, bugprone-branch-clone) #endif // _CUDA_PTX_GENERATED_MBARRIER_TEST_WAIT_H_ diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h index e13ee200b32d..ad53c6707b23 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_test_wait_parity.h @@ -218,6 +218,109 @@ _CCCL_DEVICE static inline bool mbarrier_test_wait_parity( } #endif // __cccl_ptx_isa >= 940 +/* +// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], +phaseParity; // PTX ISA 94, SM_90 +// .phase_type = { .phase_type::primary } +// .sem = { .acquire, .relaxed } +// .scope = { .cta, .cluster } +template +__device__ static inline bool mbarrier_test_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); +*/ +#if __cccl_ptx_isa >= 940 +template <::cuda::ptx::dot_sem _Sem, ::cuda::ptx::dot_scope _Scope> +_CCCL_DEVICE static inline bool mbarrier_test_wait_parity( + ::cuda::ptx::mbarrier_phase_primary_t, + ::cuda::ptx::sem_t<_Sem> __sem, + ::cuda::ptx::scope_t<_Scope> __scope, + bool& __isReportSeen, + ::cuda::std::uint8_t& __reportValue, + ::cuda::std::uint64_t* __addr, + ::cuda::std::uint32_t __phaseParity) +{ + // __phase_type == mbarrier_phase_primary (due to parameter type constraint) + static_assert(__sem == sem_acquire || __sem == sem_relaxed, ""); + static_assert(__scope == scope_cta || __scope == scope_cluster, ""); + ::cuda::std::uint32_t __waitComplete; + ::cuda::std::uint32_t __isReportSeen_tmp; + ::cuda::std::uint32_t __reportValue_tmp; + if constexpr (__sem == sem_acquire && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + else if constexpr (__sem == sem_acquire && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.test_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + __isReportSeen = static_cast(__isReportSeen_tmp); + __reportValue = static_cast<::cuda::std::uint8_t>(__reportValue_tmp); + return static_cast(__waitComplete); +} +#endif // __cccl_ptx_isa >= 940 + /* // mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90 diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h index 7119b4d2dbe4..bf629080ed0d 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait.h @@ -342,6 +342,113 @@ _CCCL_DEVICE static inline bool mbarrier_try_wait( } #endif // __cccl_ptx_isa >= 940 +/* +// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state; // PTX +ISA 94, SM_90 +// .phase_type = { .phase_type::primary } +// .sem = { .acquire, .relaxed } +// .scope = { .cta, .cluster } +template +__device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state); +*/ +#if __cccl_ptx_isa >= 940 +template <::cuda::ptx::dot_sem _Sem, ::cuda::ptx::dot_scope _Scope> +_CCCL_DEVICE static inline bool mbarrier_try_wait( + ::cuda::ptx::mbarrier_phase_primary_t, + ::cuda::ptx::sem_t<_Sem> __sem, + ::cuda::ptx::scope_t<_Scope> __scope, + bool& __isReportSeen, + ::cuda::std::uint8_t& __reportValue, + ::cuda::std::uint64_t* __addr, + ::cuda::std::uint64_t __state) +{ + // __phase_type == mbarrier_phase_primary (due to parameter type constraint) + static_assert(__sem == sem_acquire || __sem == sem_relaxed, ""); + static_assert(__scope == scope_cta || __scope == scope_cluster, ""); + ::cuda::std::uint32_t __waitComplete; + ::cuda::std::uint32_t __isReportSeen_tmp; + ::cuda::std::uint32_t __reportValue_tmp; + if constexpr (__sem == sem_acquire && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + else if constexpr (__sem == sem_acquire && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state) + : "memory"); + } + __isReportSeen = static_cast(__isReportSeen_tmp); + __reportValue = static_cast<::cuda::std::uint8_t>(__reportValue_tmp); + return static_cast(__waitComplete); +} +#endif // __cccl_ptx_isa >= 940 + /* // mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90 @@ -435,6 +542,115 @@ _CCCL_DEVICE static inline bool mbarrier_try_wait( } #endif // __cccl_ptx_isa >= 940 +/* +// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], state, +suspendTimeHint; // PTX ISA 94, SM_90 +// .phase_type = { .phase_type::primary } +// .sem = { .acquire, .relaxed } +// .scope = { .cta, .cluster } +template +__device__ static inline bool mbarrier_try_wait( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint64_t state, + uint32_t suspendTimeHint); +*/ +#if __cccl_ptx_isa >= 940 +template <::cuda::ptx::dot_sem _Sem, ::cuda::ptx::dot_scope _Scope> +_CCCL_DEVICE static inline bool mbarrier_try_wait( + ::cuda::ptx::mbarrier_phase_primary_t, + ::cuda::ptx::sem_t<_Sem> __sem, + ::cuda::ptx::scope_t<_Scope> __scope, + bool& __isReportSeen, + ::cuda::std::uint8_t& __reportValue, + ::cuda::std::uint64_t* __addr, + ::cuda::std::uint64_t __state, + ::cuda::std::uint32_t __suspendTimeHint) +{ + // __phase_type == mbarrier_phase_primary (due to parameter type constraint) + static_assert(__sem == sem_acquire || __sem == sem_relaxed, ""); + static_assert(__scope == scope_cta || __scope == scope_cluster, ""); + ::cuda::std::uint32_t __waitComplete; + ::cuda::std::uint32_t __isReportSeen_tmp; + ::cuda::std::uint32_t __reportValue_tmp; + if constexpr (__sem == sem_acquire && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state), "r"(__suspendTimeHint) + : "memory"); + } + else if constexpr (__sem == sem_acquire && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state), "r"(__suspendTimeHint) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state), "r"(__suspendTimeHint) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 P_OUT_waitComplete|P_OUT_isReportSeen, " + "B_OUT_reportValue, " + "[%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "l"(__state), "r"(__suspendTimeHint) + : "memory"); + } + __isReportSeen = static_cast(__isReportSeen_tmp); + __reportValue = static_cast<::cuda::std::uint8_t>(__reportValue_tmp); + return static_cast(__waitComplete); +} +#endif // __cccl_ptx_isa >= 940 + // NOLINTEND(modernize-unary-static-assert, bugprone-branch-clone) #endif // _CUDA_PTX_GENERATED_MBARRIER_TRY_WAIT_H_ diff --git a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h index 23846c0a22a6..7be7fa2e784f 100644 --- a/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h +++ b/libcudacxx/include/cuda/__ptx/instructions/generated/mbarrier_try_wait_parity.h @@ -348,6 +348,109 @@ _CCCL_DEVICE static inline bool mbarrier_try_wait_parity( } #endif // __cccl_ptx_isa >= 940 +/* +// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], +phaseParity; // PTX ISA 94, SM_90 +// .phase_type = { .phase_type::primary } +// .sem = { .acquire, .relaxed } +// .scope = { .cta, .cluster } +template +__device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity); +*/ +#if __cccl_ptx_isa >= 940 +template <::cuda::ptx::dot_sem _Sem, ::cuda::ptx::dot_scope _Scope> +_CCCL_DEVICE static inline bool mbarrier_try_wait_parity( + ::cuda::ptx::mbarrier_phase_primary_t, + ::cuda::ptx::sem_t<_Sem> __sem, + ::cuda::ptx::scope_t<_Scope> __scope, + bool& __isReportSeen, + ::cuda::std::uint8_t& __reportValue, + ::cuda::std::uint64_t* __addr, + ::cuda::std::uint32_t __phaseParity) +{ + // __phase_type == mbarrier_phase_primary (due to parameter type constraint) + static_assert(__sem == sem_acquire || __sem == sem_relaxed, ""); + static_assert(__scope == scope_cta || __scope == scope_cluster, ""); + ::cuda::std::uint32_t __waitComplete; + ::cuda::std::uint32_t __isReportSeen_tmp; + ::cuda::std::uint32_t __reportValue_tmp; + if constexpr (__sem == sem_acquire && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + else if constexpr (__sem == sem_acquire && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity) + : "memory"); + } + __isReportSeen = static_cast(__isReportSeen_tmp); + __reportValue = static_cast<::cuda::std::uint8_t>(__reportValue_tmp); + return static_cast(__waitComplete); +} +#endif // __cccl_ptx_isa >= 940 + /* // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90 // .phase_type = { .phase_type::conditional } @@ -515,6 +618,111 @@ _CCCL_DEVICE static inline bool mbarrier_try_wait_parity( } #endif // __cccl_ptx_isa >= 940 +/* +// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, reportValue, [addr], +phaseParity, suspendTimeHint; // PTX ISA 94, SM_90 +// .phase_type = { .phase_type::primary } +// .sem = { .acquire, .relaxed } +// .scope = { .cta, .cluster } +template +__device__ static inline bool mbarrier_try_wait_parity( + cuda::ptx::mbarrier_phase_primary_t, + cuda::ptx::sem_t sem, + cuda::ptx::scope_t scope, + bool& isReportSeen, + uint8_t& reportValue, + uint64_t* addr, + uint32_t phaseParity, + uint32_t suspendTimeHint); +*/ +#if __cccl_ptx_isa >= 940 +template <::cuda::ptx::dot_sem _Sem, ::cuda::ptx::dot_scope _Scope> +_CCCL_DEVICE static inline bool mbarrier_try_wait_parity( + ::cuda::ptx::mbarrier_phase_primary_t, + ::cuda::ptx::sem_t<_Sem> __sem, + ::cuda::ptx::scope_t<_Scope> __scope, + bool& __isReportSeen, + ::cuda::std::uint8_t& __reportValue, + ::cuda::std::uint64_t* __addr, + ::cuda::std::uint32_t __phaseParity, + ::cuda::std::uint32_t __suspendTimeHint) +{ + // __phase_type == mbarrier_phase_primary (due to parameter type constraint) + static_assert(__sem == sem_acquire || __sem == sem_relaxed, ""); + static_assert(__scope == scope_cta || __scope == scope_cluster, ""); + ::cuda::std::uint32_t __waitComplete; + ::cuda::std::uint32_t __isReportSeen_tmp; + ::cuda::std::uint32_t __reportValue_tmp; + if constexpr (__sem == sem_acquire && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity), "r"(__suspendTimeHint) + : "memory"); + } + else if constexpr (__sem == sem_acquire && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity), "r"(__suspendTimeHint) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cta) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity), "r"(__suspendTimeHint) + : "memory"); + } + else if constexpr (__sem == sem_relaxed && __scope == scope_cluster) + { + asm("{\n\t" + ".reg .pred P_OUT_waitComplete; \n\t" + ".reg .pred P_OUT_isReportSeen; \n\t" + ".reg .b8 B_OUT_reportValue; \n\t" + "mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 " + "P_OUT_waitComplete|P_OUT_isReportSeen, B_OUT_reportValue, [%3], %4, %5; \n\t" + "selp.b32 %0, 1, 0, P_OUT_waitComplete; \n\t" + "selp.b32 %1, 1, 0, P_OUT_isReportSeen; \n\t" + "cvt.u32.u8 %2, B_OUT_reportValue; \n" + "}" + : "=r"(__waitComplete), "=r"(__isReportSeen_tmp), "=r"(__reportValue_tmp) + : "r"(__as_ptr_smem(__addr)), "r"(__phaseParity), "r"(__suspendTimeHint) + : "memory"); + } + __isReportSeen = static_cast(__isReportSeen_tmp); + __reportValue = static_cast<::cuda::std::uint8_t>(__reportValue_tmp); + return static_cast(__waitComplete); +} +#endif // __cccl_ptx_isa >= 940 + /* // mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90 diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.h index e8524dbc822d..bee122551d9a 100644 --- a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.h +++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait.h @@ -105,4 +105,59 @@ __global__ void test_mbarrier_test_wait(void** fn_ptr) cuda::std::uint64_t*, cuda::std::uint64_t)>(cuda::ptx::mbarrier_test_wait));)); #endif // __cccl_ptx_isa >= 940 + +#if __cccl_ptx_isa >= 940 + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.phase_type::primary.acquire.cta.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.phase_type::primary.acquire.cluster.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.phase_type::primary.relaxed.cta.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait));)); +#endif // __cccl_ptx_isa >= 940 } diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.h index e59751977f1a..af51667b4bb9 100644 --- a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.h +++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_test_wait_parity.h @@ -108,6 +108,61 @@ __global__ void test_mbarrier_test_wait_parity(void** fn_ptr) cuda::std::uint32_t)>(cuda::ptx::mbarrier_test_wait_parity));)); #endif // __cccl_ptx_isa >= 940 +#if __cccl_ptx_isa >= 940 + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 waitComplete|isReportSeen, + // [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.test_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 waitComplete|isReportSeen, + // [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_test_wait_parity));)); +#endif // __cccl_ptx_isa >= 940 + #if __cccl_ptx_isa >= 940 NV_IF_TARGET( NV_PROVIDES_SM_90, diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.h index 195e4492f13d..a3ad7c20b451 100644 --- a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.h +++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait.h @@ -156,6 +156,61 @@ __global__ void test_mbarrier_try_wait(void** fn_ptr) cuda::std::uint64_t)>(cuda::ptx::mbarrier_try_wait));)); #endif // __cccl_ptx_isa >= 940 +#if __cccl_ptx_isa >= 940 + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); +#endif // __cccl_ptx_isa >= 940 + #if __cccl_ptx_isa >= 940 NV_IF_TARGET( NV_PROVIDES_SM_90, @@ -210,4 +265,63 @@ __global__ void test_mbarrier_try_wait(void** fn_ptr) cuda::std::uint64_t, cuda::std::uint32_t)>(cuda::ptx::mbarrier_try_wait));)); #endif // __cccl_ptx_isa >= 940 + +#if __cccl_ptx_isa >= 940 + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64 waitComplete|isReportSeen, reportValue, + // [addr], state, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait));)); +#endif // __cccl_ptx_isa >= 940 } diff --git a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.h b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.h index b7a06dd69943..bb37d37ddb7e 100644 --- a/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.h +++ b/libcudacxx/test/libcudacxx/cuda/ptx/generated/mbarrier_try_wait_parity.h @@ -160,6 +160,61 @@ __global__ void test_mbarrier_try_wait_parity(void** fn_ptr) cuda::std::uint32_t)>(cuda::ptx::mbarrier_try_wait_parity));)); #endif // __cccl_ptx_isa >= 940 +#if __cccl_ptx_isa >= 940 + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 waitComplete|isReportSeen, + // [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 waitComplete|isReportSeen, + // [addr], phaseParity; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); +#endif // __cccl_ptx_isa >= 940 + #if __cccl_ptx_isa >= 940 NV_IF_TARGET( NV_PROVIDES_SM_90, @@ -262,6 +317,65 @@ __global__ void test_mbarrier_try_wait_parity(void** fn_ptr) cuda::std::uint32_t)>(cuda::ptx::mbarrier_try_wait_parity));)); #endif // __cccl_ptx_isa >= 940 +#if __cccl_ptx_isa >= 940 + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], phaseParity, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64 waitComplete|isReportSeen, + // [addr], phaseParity, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64 waitComplete|isReportSeen, + // reportValue, [addr], phaseParity, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); + NV_IF_TARGET( + NV_PROVIDES_SM_90, + ( + // mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64 waitComplete|isReportSeen, + // [addr], phaseParity, suspendTimeHint; + * fn_ptr++ = reinterpret_cast( + static_cast(cuda::ptx::mbarrier_try_wait_parity));)); +#endif // __cccl_ptx_isa >= 940 + #if __cccl_ptx_isa >= 940 NV_IF_TARGET( NV_PROVIDES_SM_90,