tcgen05.alloc
PTX ISA: tcgen05.alloc
tcgen05.dealloc.cta_group::1.sync.aligned.b32
// tcgen05.dealloc.cta_group.sync.aligned.b32 taddr, nCols; // 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_dealloc(
cuda::ptx::cta_group_t<Cta_Group> cta_group,
uint32_t taddr,
const uint32_t& nCols);
tcgen05.dealloc.cta_group::2.sync.aligned.b32
// tcgen05.dealloc.cta_group.sync.aligned.b32 taddr, nCols; // 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_dealloc(
cuda::ptx::cta_group_t<Cta_Group> cta_group,
uint32_t taddr,
const uint32_t& nCols);
tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned
// tcgen05.relinquish_alloc_permit.cta_group.sync.aligned; // 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_relinquish_alloc_permit(
cuda::ptx::cta_group_t<Cta_Group> cta_group);
tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned
// tcgen05.relinquish_alloc_permit.cta_group.sync.aligned; // 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_relinquish_alloc_permit(
cuda::ptx::cta_group_t<Cta_Group> cta_group);