cp.async.bulk#

Implementation notes#

NOTE. Both srcMem and dstMem must be 16-byte aligned, and size must be a multiple of 16.

Changelog#

  • In earlier versions, cp_async_bulk_multicast was enabled for SM_90. This has been changed to SM_90a.

Unicast#

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 80, SM_90
// .dst       = { .shared::cluster }
// .src       = { .global }
template <typename = void>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 86, SM_90
// .dst       = { .shared::cta }
// .src       = { .global }
template <typename = void>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_shared_t,
  cuda::ptx::space_global_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cluster.shared::cta.mbarrier::complete_tx::bytes#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes [dstMem], [srcMem], size, [rdsmem_bar]; // PTX ISA 80, SM_90
// .dst       = { .shared::cluster }
// .src       = { .shared::cta }
template <typename = void>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_shared_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* rdsmem_bar);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.ignore_oob#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.ignore_oob [dstMem], [srcMem], size, ignoreBytesLeft, ignoreBytesRight, [smem_bar]; // PTX ISA 92, SM_90
// .dst       = { .shared::cta }
// .src       = { .global }
template <typename = void>
__device__ static inline void cp_async_bulk_ignore_oob(
  cuda::ptx::space_shared_t,
  cuda::ptx::space_global_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  const uint32_t& ignoreBytesLeft,
  const uint32_t& ignoreBytesRight,
  uint64_t* smem_bar);

cp.async.bulk.global.shared::cta.bulk_group#

// cp.async.bulk.dst.src.bulk_group [dstMem], [srcMem], size; // PTX ISA 80, SM_90
// .dst       = { .global }
// .src       = { .shared::cta }
template <typename = void>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_global_t,
  cuda::ptx::space_shared_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size);

cp.async.bulk.global.shared::cta.bulk_group.cp_mask#

// cp.async.bulk.dst.src.bulk_group.cp_mask [dstMem], [srcMem], size, byteMask; // PTX ISA 86, SM_100
// .dst       = { .global }
// .src       = { .shared::cta }
template <typename = void>
__device__ static inline void cp_async_bulk_cp_mask(
  cuda::ptx::space_global_t,
  cuda::ptx::space_shared_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  const uint16_t& byteMask);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80000000#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cta.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_element::ff#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80000000#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8000#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::80#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_16bytes::8#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.mbarrier::report::validity::per_element::ff#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.report_mechanism [dstMem], [srcMem], size, [smem_bar]; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cta, .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_space Space, cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_t<Space> space,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar);

Multicast#

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 80, SM_90a, SM_100a, SM_100f, SM_103a, SM_103f, SM_107a, SM_107f, SM_110a, SM_110f
// .dst       = { .shared::cluster }
// .src       = { .global }
template <typename = void>
__device__ static inline void cp_async_bulk(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint16_t& ctaMask);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster::32b [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cluster }
// .src       = { .global }
template <typename = void>
__device__ static inline void cp_async_bulk_multicast_32b(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint32_t& ctaMask);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.mbarrier::report::validity::per_16bytes::80000000#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster::32b.report_mechanism [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk_multicast_32b(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint32_t& ctaMask);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.mbarrier::report::validity::per_16bytes::8000#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster::32b.report_mechanism [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk_multicast_32b(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint32_t& ctaMask);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.mbarrier::report::validity::per_16bytes::80#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster::32b.report_mechanism [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk_multicast_32b(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint32_t& ctaMask);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.mbarrier::report::validity::per_16bytes::8#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster::32b.report_mechanism [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk_multicast_32b(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint32_t& ctaMask);

cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster::32b.mbarrier::report::validity::per_element::ff#

// cp.async.bulk.dst.src.mbarrier::complete_tx::bytes.multicast::cluster::32b.report_mechanism [dstMem], [srcMem], size, [smem_bar], ctaMask; // PTX ISA 94, SM_107a, SM_107f
// .dst       = { .shared::cluster }
// .src       = { .global }
// .report_mechanism = { .mbarrier::report::validity::per_16bytes::80000000, .mbarrier::report::validity::per_16bytes::8000, .mbarrier::report::validity::per_16bytes::80, .mbarrier::report::validity::per_16bytes::8, .mbarrier::report::validity::per_element::ff }
template <cuda::ptx::dot_report_mechanism Report_Mechanism>
__device__ static inline void cp_async_bulk_multicast_32b(
  cuda::ptx::space_cluster_t,
  cuda::ptx::space_global_t,
  cuda::ptx::report_mechanism_t<Report_Mechanism> report_mechanism,
  void* dstMem,
  const void* srcMem,
  const uint32_t& size,
  uint64_t* smem_bar,
  const uint32_t& ctaMask);
On this page