tcgen05.commit

PTX ISA: tcgen05.commit

tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64

// tcgen05.commit.cta_group.mbarrier::arrive::one.shared::cluster.b64 [smem_bar]; // PTX ISA 86, SM_100a, SM_101a
// .cta_group = { .cta_group::1, .cta_group::2 }
template <cuda::ptx::dot_cta_group Cta_Group>
__device__ static inline void tcgen05_commit(
  cuda::ptx::cta_group_t<Cta_Group> cta_group,
  uint64_t* smem_bar);

tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.b64

// tcgen05.commit.cta_group.mbarrier::arrive::one.shared::cluster.b64 [smem_bar]; // PTX ISA 86, SM_100a, SM_101a
// .cta_group = { .cta_group::1, .cta_group::2 }
template <cuda::ptx::dot_cta_group Cta_Group>
__device__ static inline void tcgen05_commit(
  cuda::ptx::cta_group_t<Cta_Group> cta_group,
  uint64_t* smem_bar);

tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64

// tcgen05.commit.cta_group.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [smem_bar], ctaMask; // PTX ISA 86, SM_100a, SM_101a
// .cta_group = { .cta_group::1, .cta_group::2 }
template <cuda::ptx::dot_cta_group Cta_Group>
__device__ static inline void tcgen05_commit_multicast(
  cuda::ptx::cta_group_t<Cta_Group> cta_group,
  uint64_t* smem_bar,
  uint16_t ctaMask);

tcgen05.commit.cta_group::2.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64

// tcgen05.commit.cta_group.mbarrier::arrive::one.shared::cluster.multicast::cluster.b64 [smem_bar], ctaMask; // PTX ISA 86, SM_100a, SM_101a
// .cta_group = { .cta_group::1, .cta_group::2 }
template <cuda::ptx::dot_cta_group Cta_Group>
__device__ static inline void tcgen05_commit_multicast(
  cuda::ptx::cta_group_t<Cta_Group> cta_group,
  uint64_t* smem_bar,
  uint16_t ctaMask);