ForwardMeta is used to store the global meta information of the model forward.
| 61 | |
| 62 | @dataclass |
| 63 | class ForwardMeta: |
| 64 | """ |
| 65 | ForwardMeta is used to store the global meta information of the model forward. |
| 66 | """ |
| 67 | |
| 68 | # Input tokens IDs of removed padding |
| 69 | ids_remove_padding: paddle.Tensor |
| 70 | # Rotation position embedding |
| 71 | rotary_embs: Optional[paddle.Tensor] = None |
| 72 | |
| 73 | # Use cuda graph in this step or not. Used to avoid run cuda graph when in dummy run or prefill stage. |
| 74 | step_use_cudagraph: bool = False |
| 75 | |
| 76 | # Flag indicating RoPE was already applied externally (e.g., by PaddleFormers) |
| 77 | # When True, FlashAttentionBackend uses identity RoPE (cos=1, sin=0) to avoid double application |
| 78 | rope_already_applied: bool = False |
| 79 | |
| 80 | # Attention backend object |
| 81 | attn_backend: AttentionBackend = None |
| 82 | # Forward mode used during attention |
| 83 | forward_mode: ForwardMode = ForwardMode.MIXED |
| 84 | # Attention mask |
| 85 | attn_mask: Optional[paddle.Tensor] = None |
| 86 | # Attention mask offset |
| 87 | attn_mask_offsets: Optional[paddle.Tensor] = None |
| 88 | |
| 89 | # A common pattern for launching CUDA kernels is to set the kernel's grids.x dimension |
| 90 | # using a `num_blocks` variable, and then map each thread block to a specific batch and |
| 91 | # data tile using `batch_ids` and `tile_ids_per_batch`. |
| 92 | # |
| 93 | # The variable names below follow this pattern, using a common prefix (e.g., `encoder_`, `decoder_`, `kv_`) |
| 94 | # for variables that are logically grouped together. The mapping works as follows: |
| 95 | # |
| 96 | # Usage: `my_kernel<<<grids, ...>>>(..., batch_ids, tile_ids, ...)` |
| 97 | # `grids.x` = `num_blocks_cpu` |
| 98 | # `batch_id` = `batch_ids[blockIdx.x]` |
| 99 | # `tile_id` = `tile_ids[blockIdx.x]` |
| 100 | |
| 101 | # Maps the thread block index (blockIdx.x) to the corresponding batch for the decoder stage in multi_query_append_attention_warp1_4_kernel. |
| 102 | # Decoder batch id. Used by attention backend. |
| 103 | decoder_batch_ids: Optional[paddle.Tensor] = None |
| 104 | # Maps the thread block index (blockIdx.x) to the specific data tile being processed within that batch for the decoder stage in multi_query_append_attention_warp1_4_kernel. |
| 105 | decoder_tile_ids_per_batch: Optional[paddle.Tensor] = None |
| 106 | # The number of blocks that attention backend can use in decode stage |
| 107 | decoder_num_blocks_device: Optional[paddle.Tensor] = None |
| 108 | # The number of CUDA blocks to launch in the x-dimension for the multi_query_append_attention_warp1_4_kernel, defining its grids.x. |
| 109 | decoder_num_blocks_cpu: Optional[paddle.Tensor] = None |
| 110 | # A tensor that holds multiple lengths related to prefill or decode stages. |
| 111 | max_len_tensor_cpu: Optional[paddle.Tensor] = None |
| 112 | # Maps the thread block index (blockIdx.x) to the corresponding batch for the encoder stage in multi_query_append_attention_kernel. |
| 113 | encoder_batch_ids: Optional[paddle.Tensor] = None |
| 114 | # Maps the thread block index (blockIdx.x) to the specific data tile being processed within that batch for the encoder stage in multi_query_append_attention_kernel. |
| 115 | encoder_tile_ids_per_batch: Optional[paddle.Tensor] = None |
| 116 | # The number of CUDA blocks to launch in the x-dimension for the multi_query_append_attention_kernel, defining its grids.x. |
| 117 | encoder_num_blocks_x_cpu: Optional[paddle.Tensor] = None |
| 118 | # Maps the thread block index (blockIdx.x) to the corresponding batch for the append_write_cache_kv kernel. |
| 119 | kv_batch_ids: Optional[paddle.Tensor] = None |
| 120 | # Maps the thread block index (blockIdx.x) to the specific data tile being processed within that batch for the append_write_cache_kv kernel. |
no outgoing calls