(self)
| 124 | """The CpAsync instruction submodule.""" |
| 125 | |
| 126 | def __init__(self): |
| 127 | self.commit_group = _op_wrapper(_cuda_op.ptx_cp_async_commit_group) |
| 128 | self.wait_group = _op_wrapper(_cuda_op.ptx_cp_async_wait_group) |
| 129 | # Legacy variant: takes (dst_ptr, dst_offset, src_ptr, src_offset, |
| 130 | # cp_size). Offsets are folded into the pointers; coexists with |
| 131 | # the fork-native ``__call__`` form. |
| 132 | self.legacy = _dtype_forward(_cuda_op.ptx_cp_async_legacy) |
| 133 | self.bulk = CpAsyncBulkNamespace() |
| 134 | self.mbarrier = CpAsyncMbarrierNamespace() |
| 135 | |
| 136 | def __call__(self, *args, **kwds): |
| 137 | # Accept the legacy 6-arg form ``(elem_dtype, dst, dst_off, src, |
nothing calls this directly
no test coverage detected