(self)
| 42 | """The PTX instruction submodule.""" |
| 43 | |
| 44 | def __init__(self): |
| 45 | self.ldg32 = _ptx_ldg32 |
| 46 | self.ldmatrix = _dtype_forward(_cuda_op.ptx_ldmatrix) |
| 47 | # Apache-compatible variant. Same lowered intrinsic as |
| 48 | # ``ldmatrix`` but accepts the historical ``(trans, num, dtype, |
| 49 | # local_ptr, local_offset, smem_ptr, smem_offset)`` form. Coexists |
| 50 | # with the fork-native version so upstream-derived tests keep |
| 51 | # working without rewriting their tirx code. |
| 52 | self.ldmatrix_legacy = _dtype_forward(_cuda_op.ptx_ldmatrix_legacy) |
| 53 | self.stmatrix = _op_wrapper(_cuda_op.ptx_stmatrix) |
| 54 | self.setmaxnreg: Callable[..., Any] = _op_wrapper(_cuda_op.ptx_setmaxnreg) |
| 55 | self.elect_sync: Callable[..., Any] = _op_wrapper(_cuda_op.ptx_elect_sync) |
| 56 | self.clc_try_cancel = _op_wrapper(_cuda_op.ptx_clc_try_cancel) |
| 57 | self.clc_query_cancel = _op_wrapper(_cuda_op.ptx_clc_query_cancel) |
| 58 | self.fetch_register: Callable[..., Any] = _op_wrapper(_cuda_op.ptx_fetch_register) |
| 59 | self.ld = _op_wrapper(_cuda_op.ptx_ld) |
| 60 | self.ld_acquire = _op_wrapper(_cuda_op.ptx_ld_acquire) |
| 61 | self.ld_volatile = _op_wrapper(_cuda_op.ptx_ld_volatile) |
| 62 | self.ld_global_acquire = _op_wrapper(_cuda_op.ptx_ld_global_acquire) |
| 63 | self.red_scalar = _op_wrapper(_cuda_op.ptx_red_scalar) |
| 64 | self.atom_scalar = _op_wrapper(_cuda_op.ptx_atom_scalar) |
| 65 | self.prefetch_tensormap = _op_wrapper(_cuda_op.ptx_prefetch_tensormap) |
| 66 | self.mbarrier_test_wait_parity = _op_wrapper(_cuda_op.ptx_mbarrier_test_wait_parity) |
| 67 | self.cp_async_bulk_g2s_cta = _op_wrapper(_cuda_op.ptx_cp_async_bulk_g2s_cta) |
| 68 | self.cp_async_bulk_g2s_cluster = _op_wrapper(_cuda_op.ptx_cp_async_bulk_g2s_cluster) |
| 69 | self.cp_async_bulk_s2s_cluster = _op_wrapper(_cuda_op.ptx_cp_async_bulk_s2s_cluster) |
| 70 | self.cp_async_bulk_s2g = _op_wrapper(_cuda_op.ptx_cp_async_bulk_s2g) |
| 71 | self.st = _op_wrapper(_cuda_op.ptx_st) |
| 72 | self.st_bulk = _op_wrapper(_cuda_op.ptx_st_bulk) |
| 73 | self.fns_b32 = _op_wrapper(_cuda_op.ptx_fns_b32) |
| 74 | self.add_rn_f32_bf16 = _op_wrapper(_cuda_op.ptx_add_rn_f32_bf16) |
| 75 | self.mapa = _op_wrapper(_cuda_op.ptx_mapa) |
| 76 | self.map_shared_rank = _op_wrapper(_cuda_op.ptx_map_shared_rank) |
| 77 | self.any_sync = _op_wrapper(_cuda_op.ptx_any_sync) |
| 78 | # Math operations |
| 79 | self.exp2 = _op_wrapper(_cuda_op.ptx_exp2) |
| 80 | self.rcp = _op_wrapper(_cuda_op.ptx_rcp) |
| 81 | self.reduce3_min_f32 = _op_wrapper(_cuda_op.ptx_reduce3_min_f32) |
| 82 | self.reduce3_max_f32 = _op_wrapper(_cuda_op.ptx_reduce3_max_f32) |
| 83 | # add/sub/mul/fma DPS form: (d_addr, a, b[, c], *, rounding, ftz[, sat]) |
| 84 | self.add_f32 = _op_wrapper(_cuda_op.ptx_add_f32) |
| 85 | self.add_f32x2 = _op_wrapper(_cuda_op.ptx_add_f32x2) |
| 86 | self.add_f64 = _op_wrapper(_cuda_op.ptx_add_f64) |
| 87 | self.sub_f32 = _op_wrapper(_cuda_op.ptx_sub_f32) |
| 88 | self.sub_f32x2 = _op_wrapper(_cuda_op.ptx_sub_f32x2) |
| 89 | self.sub_f64 = _op_wrapper(_cuda_op.ptx_sub_f64) |
| 90 | self.mul_f32 = _op_wrapper(_cuda_op.ptx_mul_f32) |
| 91 | self.mul_f32x2 = _op_wrapper(_cuda_op.ptx_mul_f32x2) |
| 92 | self.mul_f64 = _op_wrapper(_cuda_op.ptx_mul_f64) |
| 93 | self.fma_f32 = _op_wrapper(_cuda_op.ptx_fma_f32) |
| 94 | self.fma_f32x2 = _op_wrapper(_cuda_op.ptx_fma_f32x2) |
| 95 | self.fma_f64 = _op_wrapper(_cuda_op.ptx_fma_f64) |
| 96 | self.max_f32 = _op_wrapper(_cuda_op.ptx_max_f32) |
| 97 | self.mma = MmaNamespace() |
| 98 | self.cp_async = CpAsyncNamespace() |
| 99 | self.wgmma = WgmmaNamespace() |
| 100 | self.mbarrier = MbarrierNamespace() |
| 101 | self.tcgen05 = Tcgen05Namespace() |
nothing calls this directly
no test coverage detected