mbarrier.test_wait / mbarrier.try_wait

• PTX ISA: mbarrier.test_wait / mbarrier.try_wait

mbarrier.test_wait

mbarrier.test_wait.shared.b64

// mbarrier.test_wait.shared.b64 waitComplete, [addr], state; // PTX ISA 70, SM_80
template <typename = void>
__device__ static inline bool mbarrier_test_wait(
  uint64_t* addr,
  const uint64_t& state);

mbarrier.test_wait.acquire.cta.shared::cta.b64

// mbarrier.test_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.test_wait.acquire.cluster.shared::cta.b64

// mbarrier.test_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.test_wait.relaxed.cta.shared::cta.b64

// mbarrier.test_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.test_wait.relaxed.cluster.shared::cta.b64

// mbarrier.test_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.test_wait.phase_type::primary.acquire.cta.shared::cta.b64

// mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.acquire.cluster.shared::cta.b64

// mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.relaxed.cta.shared::cta.b64

// mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.relaxed.cluster.shared::cta.b64

// mbarrier.test_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.acquire.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.acquire.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.relaxed.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.phase_type::primary.relaxed.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.test_wait.parity

mbarrier.test_wait.parity.shared.b64

// mbarrier.test_wait.parity.shared.b64 waitComplete, [addr], phaseParity; // PTX ISA 71, SM_80
template <typename = void>
__device__ static inline bool mbarrier_test_wait_parity(
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.test_wait.parity.acquire.cta.shared::cta.b64

// mbarrier.test_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.test_wait.parity.acquire.cluster.shared::cta.b64

// mbarrier.test_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.test_wait.parity.relaxed.cta.shared::cta.b64

// mbarrier.test_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.test_wait.parity.relaxed.cluster.shared::cta.b64

// mbarrier.test_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.test_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::conditional.acquire.cta.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::conditional.acquire.cluster.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::conditional.relaxed.cta.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.test_wait.parity.phase_type::conditional.relaxed.cluster.shared::cta.b64

// mbarrier.test_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_test_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait

mbarrier.try_wait.shared::cta.b64

// mbarrier.try_wait.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline bool mbarrier_try_wait(
  uint64_t* addr,
  const uint64_t& state);

mbarrier.try_wait.shared::cta.b64

// mbarrier.try_wait.shared::cta.b64 waitComplete, [addr], state, suspendTimeHint; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline bool mbarrier_try_wait(
  uint64_t* addr,
  const uint64_t& state,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.acquire.cta.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.try_wait.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.try_wait.acquire.cta.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state, suspendTimeHint; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state, suspendTimeHint; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.try_wait.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state);

mbarrier.try_wait.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state, suspendTimeHint; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.sem.scope.shared::cta.b64 waitComplete, [addr], state, suspendTimeHint; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint64_t& state,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state);

mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], state, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.acquire.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.acquire.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.relaxed.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.phase_type::primary.relaxed.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint64_t state,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity

mbarrier.try_wait.parity.shared::cta.b64

// mbarrier.try_wait.parity.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline bool mbarrier_try_wait_parity(
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.try_wait.parity.shared::cta.b64

// mbarrier.try_wait.parity.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 78, SM_90
template <typename = void>
__device__ static inline bool mbarrier_try_wait_parity(
  uint64_t* addr,
  const uint32_t& phaseParity,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.parity.acquire.cta.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.try_wait.parity.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.try_wait.parity.acquire.cta.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.parity.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 80, SM_90
// .sem       = { .acquire }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_acquire_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.parity.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.try_wait.parity.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity);

mbarrier.try_wait.parity.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.parity.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.parity.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 86, SM_90
// .sem       = { .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::sem_relaxed_t,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  const uint32_t& phaseParity,
  const uint32_t& suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::conditional.acquire.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::conditional.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::conditional.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::conditional.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity);

mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete|isReportSeen, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::primary }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.acquire.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.acquire.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cta.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::primary.relaxed.cluster.shared::cta.b64

// 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 <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__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,
  uint8_t& reportValue,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::conditional.acquire.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::conditional.acquire.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::conditional.relaxed.cta.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);

mbarrier.try_wait.parity.phase_type::conditional.relaxed.cluster.shared::cta.b64

// mbarrier.try_wait.parity.phase_type.sem.scope.shared::cta.b64 waitComplete, [addr], phaseParity, suspendTimeHint; // PTX ISA 94, SM_90
// .phase_type = { .phase_type::conditional }
// .sem       = { .acquire, .relaxed }
// .scope     = { .cta, .cluster }
template <cuda::ptx::dot_sem Sem, cuda::ptx::dot_scope Scope>
__device__ static inline bool mbarrier_try_wait_parity(
  cuda::ptx::mbarrier_phase_conditional_t,
  cuda::ptx::sem_t<Sem> sem,
  cuda::ptx::scope_t<Scope> scope,
  uint64_t* addr,
  uint32_t phaseParity,
  uint32_t suspendTimeHint);