--- name: fuse-triton-ascend-cat-conv1d description: Refactor fused causal Conv1d state updates for Triton-Ascend by replacing negative-offset cat emulation and tl.where selection with transposed UB assembly through extension.insert_slice. Use when an agent ports causal_conv1d_update-style kernels, sees discrete scalar loads from negative offsets, suffers 32-byte tail-axis padding, or launches far more logical tasks than Ascend vector cores. --- # Fuse Triton-Ascend Cat and Causal Conv1d ## Goal Join convolution state and new tokens in UB, update the cached tail, and compute grouped causal Conv1d without materializing a global concatenation. ## Workflow 1. Preserve the reference contract for x, conv_state, depthwise weights, bias, activation, and state-row indices. 2. Load state and input from nonnegative contiguous addresses. 3. Transpose or reshape them so concatenation occurs along a contiguous flattened UB axis. 4. Insert state followed by new tokens with extension.insert_slice. 5. Derive the new state tail and convolution windows from the assembled tensor. 6. Map work to approximately the available vector-core count and loop over remaining feature tiles inside each program. ## Implementation Pattern ~~~python import triton.language.extra.cann.extension as extension state_t = load_state_as_contiguous_transposed_tile() x_t = load_input_as_contiguous_transposed_tile() joined = tl.zeros((CAT_LEN * DIM_BLOCK,), dtype=x_ptr.dtype.element_ty) joined = extension.insert_slice( joined, state_t, (0,), (STATE_LEN * DIM_BLOCK,), (1,) ) joined = extension.insert_slice( joined, x_t, (STATE_LEN * DIM_BLOCK,), (SEQ_LEN * DIM_BLOCK,), (1,) ) ~~~ ## Ascend Guardrails - Do not emulate cat with negative pointer offsets plus tl.where; Triton-Ascend may classify it as discrete scalar access. - Use extension.insert_slice, not deprecated tl.insert_slice. - Account for Ascend UB padding of the final axis to 32 bytes. - For short final axes such as 1 or 3, consider a proven reshape-transpose “borrowed axis” layout only when total storage alignment permits it. - Include transpose cost and padding in the UB budget. - Query the target vector-core count instead of hardcoding a 40/48-core assumption. - Verify state updates when state indices repeat; concurrent writes need an explicit ordering policy. ## Verification Compare output and updated conv_state with the PyTorch reference for activation=None and SiLU, several state/sequence lengths, repeated and unique state indices, short widths, and feature dimensions that trigger alignment padding.