diff --git a/TraceLens/PerfModel/extensions/attention_perf_model_extensions.py b/TraceLens/PerfModel/extensions/attention_perf_model_extensions.py index b8f8c56b7..3cac4b664 100644 --- a/TraceLens/PerfModel/extensions/attention_perf_model_extensions.py +++ b/TraceLens/PerfModel/extensions/attention_perf_model_extensions.py @@ -580,27 +580,79 @@ class mha_varlen_fwd(InferenceAttention): class aiter_fmha_v3_varlen_fwd(InferenceAttention): """ - Performance model for ``aiter::fmha_v3_varlen_fwd`` (inference: sglang / vLLM). + Annotation-aware perf model for ``aiter::fmha_v3_varlen_fwd`` (inference: + sglang / vLLM). Uses the same chunk statistics as :class:`InferenceAttention` (``annotation`` on the event). Sets ``d_h_v`` from the **v** tensor (``Input Dims[2]``) so MLA shapes with differing Q/K vs V head dims are modeled correctly. - Unparseable annotation yields :meth:`InferenceAttention.no_perf_param_details` - (see base class); no packed-tensor fallback. + If annotation parsing fails, this class falls back to the core shape-based + SDPA model so regular training traces still get FLOPs/bytes coverage while + annotated inference traces use the more accurate per-request sequence stats. """ + category = "SDPA_fwd" + bwd_category = None + + def __init__(self, event, arch=None, python_path=None, enable_origami=False): + self.enable_origami = enable_origami + super().__init__(event, arch, python_path) + + @staticmethod + def _core_param_details(event): + from TraceLens.PerfModel import perf_model + + params = perf_model.aiter__fmha_v3_varlen_fwd.get_param_details(event).copy() + params["_fallback_core"] = True + return params + + def _core_model(self): + from TraceLens.PerfModel import perf_model + + return perf_model.aiter__fmha_v3_varlen_fwd( + self.event, + self.arch, + self.python_path, + enable_origami=self.enable_origami, + ) + @staticmethod def get_param_details(event): params = InferenceAttention.get_param_details(event) if params.get("_no_perf"): - return params + try: + return aiter_fmha_v3_varlen_fwd._core_param_details(event) + except (ValueError, IndexError, KeyError, TypeError): + return params args = event.get("args") or {} dims = args.get("Input Dims") or [] if len(dims) > 2 and len(dims[2]) >= 1: params["d_h_v"] = dims[2][-1] return params + def flops(self): + if self.param_details.get("_fallback_core"): + return self._core_model().flops() + return super().flops() + + def bytes(self, bytes_per_element=None): + if self.param_details.get("_fallback_core"): + if bytes_per_element is None: + return self._core_model().bytes() + return self._core_model().bytes(bytes_per_element) + return super().bytes(bytes_per_element) + + def get_compute_precision(self): + if self.param_details.get("_fallback_core"): + return self._core_model().get_compute_precision() + return super().get_compute_precision() + + def get_simulation_time(self): + if self.param_details.get("_fallback_core"): + return self._core_model().get_simulation_time() + return None + class aiter_paged_attention_ragged(InferenceAttention): """ diff --git a/TraceLens/PerfModel/extensions/pseudo_ops_perf_utils.py b/TraceLens/PerfModel/extensions/pseudo_ops_perf_utils.py index 87208b71c..aaaeba045 100644 --- a/TraceLens/PerfModel/extensions/pseudo_ops_perf_utils.py +++ b/TraceLens/PerfModel/extensions/pseudo_ops_perf_utils.py @@ -53,6 +53,8 @@ def get_pseudo_op_mappings(): # Attention pseudo ops "vllm::unified_attention_with_output": attention_perf_model_extensions.vllm_unified_attention_with_output, "aiter::mha_varlen_fwd": attention_perf_model_extensions.mha_varlen_fwd, + # Prefer the annotation-aware inference extension when annotations are + # present; it falls back to the core SDPA model for plain training rows. "aiter::fmha_v3_varlen_fwd": attention_perf_model_extensions.aiter_fmha_v3_varlen_fwd, "aiter::mha_batch_prefill": attention_perf_model_extensions.aiter_mha_batch_prefill, "sglang_profiler::attention_paged_attention_ragged": attention_perf_model_extensions.aiter_paged_attention_ragged, diff --git a/TraceLens/PerfModel/perf_model.py b/TraceLens/PerfModel/perf_model.py index d7a670051..cfda502d3 100644 --- a/TraceLens/PerfModel/perf_model.py +++ b/TraceLens/PerfModel/perf_model.py @@ -3093,6 +3093,473 @@ def bytes(self, bytes_per_element=2): return self.bytes_bwd(bytes_per_element) +class aiter__fmha_v3_bwd(SDPA): + """Perf model for ``aiter::fmha_v3_bwd`` (non-varlen). Issue #590. + + Same argument layout as ``aiter::mha_bwd``: + dout[0], q[1], k[2], v[3], out[4], softmax_lse[5], + dropout_p[6], softmax_scale[7], is_causal[8], ... + Tensors are in raw (B, N, H, d_h) layout (bnhd); bhnd_idx=(0,2,1,3) extracts + B, H, N, d_h. Mirrors :class:`aiter__mha_bwd`. + """ + + category = "SDPA_bwd" + + @staticmethod + def get_param_details(event): + return aiter__mha_bwd.get_param_details(event) + + def flops(self): + return self.flops_bwd() + + def bytes(self, bytes_per_element=2): + return self.bytes_bwd(bytes_per_element) + + +def _parse_aiter_fmha_v3_varlen_fwd_args(event, tensor_offset=0): + """Shared parser for ``aiter::(wrapper_)fmha_v3_varlen_fwd``. Issue #650 / #290. + + Direct variant (``aiter::fmha_v3_varlen_fwd``, ``tensor_offset=0``): + Input Dims: q[0], k[1], v[2], cu_seqlens_q[3], cu_seqlens_k[4], scalars... + Concrete Inputs: ..., max_seqlen_q[5], max_seqlen_k[6], min_seqlen_q[7], + dropout_p[8], softmax_scale[9], logits_soft_cap[10], + is_causal[11], ... + + Wrapper variant (``aiter::wrapper_fmha_v3_varlen_fwd``, ``tensor_offset=1``): + Extra leading ``out`` tensor — all indices shift by +1. Mirrors the + precedent set by ``aiter__fmha_v3_forward`` vs ``aiter__fmha_v3_fwd``. + + Q/K/V are packed varlen tensors in (T, H, d_h) layout. Cross-attention is + supported: K/V may have different T than Q. + """ + input_dims = event["args"]["Input Dims"] + input_types = event["args"].get("Input type", []) + concrete_inputs = event["args"]["Concrete Inputs"] + + q_idx = tensor_offset + 0 + k_idx = tensor_offset + 1 + v_idx = tensor_offset + 2 + cu_q_idx = tensor_offset + 3 + cu_k_idx = tensor_offset + 4 + scalar_base = tensor_offset + 5 + max_q_idx = scalar_base + 0 + max_kv_idx = scalar_base + 1 + dropout_idx = scalar_base + 3 + causal_idx = scalar_base + 6 + + q_shape, k_shape, v_shape = ( + input_dims[q_idx], + input_dims[k_idx], + input_dims[v_idx], + ) + dtype_A_B = ( + input_types[q_idx] if q_idx < len(input_types) else None, + input_types[k_idx] if k_idx < len(input_types) else None, + ) + hnd_idx = 1, 0, 2 # (T, H, d_h) layout + sdpa_cfg = extract_sdpa_varlen_cfg(q_shape, k_shape, v_shape, hnd_idx) + B, N_Q, H_Q, N_KV, H_KV, d_h_qk, d_h_v = ( + sdpa_cfg[key] for key in ["B", "N_Q", "H_Q", "N_KV", "H_KV", "d_h_qk", "d_h_v"] + ) + + cu_q_dim = input_dims[cu_q_idx] if cu_q_idx < len(input_dims) else [] + cu_k_dim = input_dims[cu_k_idx] if cu_k_idx < len(input_dims) else [] + num_seqs_q = cu_q_dim[0] - 1 if cu_q_dim else 1 + num_seqs_kv = cu_k_dim[0] - 1 if cu_k_dim else 1 + + def _f(idx, default): + v = concrete_inputs[idx] if idx < len(concrete_inputs) else "" + if v in ("", "None"): + return default + try: + return float(v) + except (ValueError, TypeError): + return default + + max_seqlen_q = _f(max_q_idx, float(N_Q)) + max_seqlen_kv = _f(max_kv_idx, float(N_KV)) + dropout = _f(dropout_idx, 0.0) + causal = False + if causal_idx < len(concrete_inputs): + cv = concrete_inputs[causal_idx] + if cv not in ("", "None"): + causal = cv.lower() == "true" + + return { + "B": B, + "N_Q": N_Q, + "H_Q": H_Q, + "N_KV": N_KV, + "H_KV": H_KV, + "d_h_qk": d_h_qk, + "d_h_v": d_h_v, + "dropout": dropout, + "causal": causal, + "flash_impl": True, + "num_seqs_q": num_seqs_q, + "num_seqs_kv": num_seqs_kv, + "max_seqlen_q": max_seqlen_q, + "max_seqlen_kv": max_seqlen_kv, + "dtype_A_B": dtype_A_B, + } + + +def _parse_aiter_fmha_v3_varlen_bwd_args(event, tensor_offset=0): + """Shared parser for ``aiter::(wrapper_)fmha_v3_varlen_bwd``. Issue #650. + + Direct variant (``aiter::fmha_v3_varlen_bwd``, ``tensor_offset=0``): + Input Dims: dout[0], q[1], k[2], v[3], out[4], softmax_lse[5], + cu_seqlens_q[6], cu_seqlens_k[7], scalars... + Concrete Inputs: ..., max_seqlen_q[8], max_seqlen_k[9], + dropout_p[10], softmax_scale[11], is_causal[12], ... + + Wrapper variant (``aiter::wrapper_fmha_v3_varlen_bwd``, ``tensor_offset=1``): + Extra leading tensor — all indices shift by +1. + + Q/K/V are packed varlen tensors in (T, H, d_h) layout; cross-attention with + differing N_Q vs N_KV is supported. + """ + input_dims = event["args"]["Input Dims"] + input_types = event["args"].get("Input type", []) + concrete_inputs = event["args"]["Concrete Inputs"] + + q_idx = tensor_offset + 1 + k_idx = tensor_offset + 2 + v_idx = tensor_offset + 3 + cu_q_idx = tensor_offset + 6 + cu_k_idx = tensor_offset + 7 + scalar_base = tensor_offset + 8 + max_q_idx = scalar_base + 0 + max_kv_idx = scalar_base + 1 + dropout_idx = scalar_base + 2 + causal_idx = scalar_base + 4 + + q_shape, k_shape, v_shape = ( + input_dims[q_idx], + input_dims[k_idx], + input_dims[v_idx], + ) + dtype_A_B = ( + input_types[q_idx] if q_idx < len(input_types) else None, + input_types[k_idx] if k_idx < len(input_types) else None, + ) + hnd_idx = 1, 0, 2 # (T, H, d_h) layout + sdpa_cfg = extract_sdpa_varlen_cfg(q_shape, k_shape, v_shape, hnd_idx) + B, N_Q, H_Q, N_KV, H_KV, d_h_qk, d_h_v = ( + sdpa_cfg[key] for key in ["B", "N_Q", "H_Q", "N_KV", "H_KV", "d_h_qk", "d_h_v"] + ) + + cu_q_dim = input_dims[cu_q_idx] if cu_q_idx < len(input_dims) else [] + cu_k_dim = input_dims[cu_k_idx] if cu_k_idx < len(input_dims) else [] + num_seqs_q = cu_q_dim[0] - 1 if cu_q_dim else 1 + num_seqs_kv = cu_k_dim[0] - 1 if cu_k_dim else 1 + + def _f(idx, default): + v = concrete_inputs[idx] if idx < len(concrete_inputs) else "" + if v in ("", "None"): + return default + try: + return float(v) + except (ValueError, TypeError): + return default + + max_seqlen_q = _f(max_q_idx, float(N_Q)) + max_seqlen_kv = _f(max_kv_idx, float(N_KV)) + dropout = _f(dropout_idx, 0.0) + causal = False + if causal_idx < len(concrete_inputs): + cv = concrete_inputs[causal_idx] + if cv not in ("", "None"): + causal = cv.lower() == "true" + + return { + "B": B, + "N_Q": N_Q, + "H_Q": H_Q, + "N_KV": N_KV, + "H_KV": H_KV, + "d_h_qk": d_h_qk, + "d_h_v": d_h_v, + "dropout": dropout, + "causal": causal, + "flash_impl": True, + "num_seqs_q": num_seqs_q, + "num_seqs_kv": num_seqs_kv, + "max_seqlen_q": max_seqlen_q, + "max_seqlen_kv": max_seqlen_kv, + "dtype_A_B": dtype_A_B, + } + + +class aiter__fmha_v3_varlen_fwd(SDPA): + """Perf model for ``aiter::fmha_v3_varlen_fwd``. Issue #650. + + Varlen forward attention used in diffusion training (Wan 2.x) and inference + (sglang / vLLM). Q/K/V in (T, H, d_h) packed layout, sequence boundaries + in cu_seqlens. Shape-based FLOPs use ``SDPA.flops_func`` with + ``N_Q = max_seqlen_q``, ``N_KV = max_seqlen_kv`` (lower-bound estimate when + multiple varlen sequences are present). + + Replaces the annotation-only path in + ``extensions.attention_perf_model_extensions.aiter_fmha_v3_varlen_fwd`` for + training traces; that extension class continues to exist for sglang/vLLM + inference flows that inject chunk annotations. + """ + + def __init__(self, event, arch=None, python_path=None, enable_origami=False): + super().__init__(event, arch, python_path, enable_origami=enable_origami) + self.num_seqs_q = self.param_details["num_seqs_q"] + self.num_seqs_kv = self.param_details["num_seqs_kv"] + self.max_seqlen_q = self.param_details["max_seqlen_q"] + self.max_seqlen_kv = self.param_details["max_seqlen_kv"] + + @staticmethod + def get_param_details(event): + return _parse_aiter_fmha_v3_varlen_fwd_args(event, tensor_offset=0) + + def flops(self): + # Varlen FLOPs accumulator (mirrors ``flash_attention_varlen_forward``): + # the B and S dimensions are collapsed into a single packed T axis with + # ``num_seqs_q`` variable-length sequences whose exact per-sequence + # lengths are not in the trace. We estimate as one ``max_seqlen`` + # sequence plus ``num_seqs - 1`` equal-length sequences whose lengths + # average to ``(N - max_seqlen) / (num_seqs - 1)`` (lower bound: the + # max-length sequence dominates the N^2 term). + accum_flops = self.flops_func( + self.B, + self.max_seqlen_q, + self.H_Q, + self.max_seqlen_kv, + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + ) + if self.num_seqs_q > 1: + accum_flops += (self.num_seqs_q - 1) * self.flops_func( + self.B, + (self.N_Q - self.max_seqlen_q) // (self.num_seqs_q - 1), + self.H_Q, + (self.N_KV - self.max_seqlen_kv) // (self.num_seqs_kv - 1), + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + ) + return accum_flops + + +class aiter__fmha_v3_varlen_forward(SDPA): + """Perf model for ``aiter::wrapper_fmha_v3_varlen_fwd``. Issue #290. + + Same kernel as :class:`aiter__fmha_v3_varlen_fwd` accessed via the Python + wrapper layer; the wrapper traces an extra leading ``out`` tensor, so all + arg indices shift by +1. + """ + + def __init__(self, event, arch=None, python_path=None, enable_origami=False): + super().__init__(event, arch, python_path, enable_origami=enable_origami) + self.num_seqs_q = self.param_details["num_seqs_q"] + self.num_seqs_kv = self.param_details["num_seqs_kv"] + self.max_seqlen_q = self.param_details["max_seqlen_q"] + self.max_seqlen_kv = self.param_details["max_seqlen_kv"] + + @staticmethod + def get_param_details(event): + return _parse_aiter_fmha_v3_varlen_fwd_args(event, tensor_offset=1) + + def flops(self): + # See note on ``aiter__fmha_v3_varlen_fwd.flops()``. + accum_flops = self.flops_func( + self.B, + self.max_seqlen_q, + self.H_Q, + self.max_seqlen_kv, + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + ) + if self.num_seqs_q > 1: + accum_flops += (self.num_seqs_q - 1) * self.flops_func( + self.B, + (self.N_Q - self.max_seqlen_q) // (self.num_seqs_q - 1), + self.H_Q, + (self.N_KV - self.max_seqlen_kv) // (self.num_seqs_kv - 1), + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + ) + return accum_flops + + +class aiter__fmha_v3_varlen_bwd(SDPA): + """Perf model for ``aiter::fmha_v3_varlen_bwd``. Issue #650. + + Varlen backward attention. Q/K/V in (T, H, d_h) packed layout, dout at + Input Dims[0], cu_seqlens at [6]/[7]. FLOPs use ``SDPA.flops_bwd_func`` with + ``flash_impl=True`` (recompute QK + 4 GEMMs ⇒ 5/2 × forward FLOPs for the + square N_Q=N_KV case). Cross-attention (N_Q ≠ N_KV) supported. + """ + + category = "SDPA_bwd" + + def __init__(self, event, arch=None, python_path=None, enable_origami=False): + super().__init__(event, arch, python_path, enable_origami=enable_origami) + self.num_seqs_q = self.param_details["num_seqs_q"] + self.num_seqs_kv = self.param_details["num_seqs_kv"] + self.max_seqlen_q = self.param_details["max_seqlen_q"] + self.max_seqlen_kv = self.param_details["max_seqlen_kv"] + + @staticmethod + def get_param_details(event): + return _parse_aiter_fmha_v3_varlen_bwd_args(event, tensor_offset=0) + + def flops(self): + # See note on ``aiter__fmha_v3_varlen_fwd.flops()``; the bwd FLOPs + # identity ``bwd = 5/2 * fwd`` (square self-attn, flash_impl=True) is + # baked into ``flops_bwd_func`` so the multi-seq accumulator stays + # consistent with the fwd one. + accum_flops = self.flops_bwd_func( + self.B, + self.max_seqlen_q, + self.H_Q, + self.max_seqlen_kv, + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + self.param_details["flash_impl"], + ) + if self.num_seqs_q > 1: + accum_flops += (self.num_seqs_q - 1) * self.flops_bwd_func( + self.B, + (self.N_Q - self.max_seqlen_q) // (self.num_seqs_q - 1), + self.H_Q, + (self.N_KV - self.max_seqlen_kv) // (self.num_seqs_kv - 1), + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + self.param_details["flash_impl"], + ) + return accum_flops + + def bytes(self, bytes_per_element=2): + return self.bytes_bwd(bytes_per_element) + + +class aiter__fmha_v3_varlen_backward(SDPA): + """Perf model for ``aiter::wrapper_fmha_v3_varlen_bwd``. Issue #650. + + Same kernel as :class:`aiter__fmha_v3_varlen_bwd` accessed via the Python + wrapper layer; tensor indices shift by +1 for the extra leading tensor. + """ + + category = "SDPA_bwd" + + def __init__(self, event, arch=None, python_path=None, enable_origami=False): + super().__init__(event, arch, python_path, enable_origami=enable_origami) + self.num_seqs_q = self.param_details["num_seqs_q"] + self.num_seqs_kv = self.param_details["num_seqs_kv"] + self.max_seqlen_q = self.param_details["max_seqlen_q"] + self.max_seqlen_kv = self.param_details["max_seqlen_kv"] + + @staticmethod + def get_param_details(event): + return _parse_aiter_fmha_v3_varlen_bwd_args(event, tensor_offset=1) + + def flops(self): + # See note on ``aiter__fmha_v3_varlen_bwd.flops()``. + accum_flops = self.flops_bwd_func( + self.B, + self.max_seqlen_q, + self.H_Q, + self.max_seqlen_kv, + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + self.param_details["flash_impl"], + ) + if self.num_seqs_q > 1: + accum_flops += (self.num_seqs_q - 1) * self.flops_bwd_func( + self.B, + (self.N_Q - self.max_seqlen_q) // (self.num_seqs_q - 1), + self.H_Q, + (self.N_KV - self.max_seqlen_kv) // (self.num_seqs_kv - 1), + self.H_KV, + self.d_h_qk, + self.d_h_v, + self.param_details["causal"], + self.param_details["flash_impl"], + ) + return accum_flops + + def bytes(self, bytes_per_element=2): + return self.bytes_bwd(bytes_per_element) + + +class aten___flash_attention_forward(SDPA): + """Perf model for ``aten::_flash_attention_forward``. Issue #650. + + PyTorch's dispatcher-level Flash-Attention forward (called by + ``torch.nn.functional.scaled_dot_product_attention`` when the flash backend + is selected). Mirrors :class:`aten__scaled_dot_product_flash_attention`. + + Signature (Input Dims / Concrete Inputs are parallel arrays): + q[0], k[1], v[2], cum_seq_q[3], cum_seq_k[4], + max_q[5], max_k[6], dropout_p[7], is_causal[8], + return_debug_mask[9], scale[10], ... + Q/K/V layout: (B, S, H, d_h); bhnd_idx=(0, 2, 1, 3) extracts B, H, N, d_h. + """ + + @staticmethod + def get_param_details(event): + input_dims = event["args"]["Input Dims"] + input_types = event["args"].get("Input type", []) + concrete_inputs = event["args"]["Concrete Inputs"] + q_idx, k_idx, v_idx = 0, 1, 2 + q_shape, k_shape, v_shape = ( + input_dims[q_idx], + input_dims[k_idx], + input_dims[v_idx], + ) + dtype_A_B = ( + input_types[q_idx] if q_idx < len(input_types) else None, + input_types[k_idx] if k_idx < len(input_types) else None, + ) + bhnd_idx = 0, 2, 1, 3 + sdpa_cfg = extract_sdpa_cfg(q_shape, k_shape, v_shape, bhnd_idx) + B, N_Q, H_Q, N_KV, H_KV, d_h_qk, d_h_v = ( + sdpa_cfg[key] + for key in ["B", "N_Q", "H_Q", "N_KV", "H_KV", "d_h_qk", "d_h_v"] + ) + dropout_p = 0.0 + if len(concrete_inputs) > 7 and concrete_inputs[7] not in ("", "None"): + try: + dropout_p = float(concrete_inputs[7]) + except (ValueError, TypeError): + pass + is_causal = False + if len(concrete_inputs) > 8 and concrete_inputs[8] not in ("", "None"): + is_causal = concrete_inputs[8].lower() == "true" + return { + "B": B, + "N_Q": N_Q, + "H_Q": H_Q, + "N_KV": N_KV, + "H_KV": H_KV, + "d_h_qk": d_h_qk, + "d_h_v": d_h_v, + "dropout": dropout_p, + "causal": is_causal, + "flash_impl": True, + "dtype_A_B": dtype_A_B, + } + + class vllm_unified_attention_with_output(SDPA): @staticmethod diff --git a/TraceLens/PerfModel/torch_op_mapping.py b/TraceLens/PerfModel/torch_op_mapping.py index a421bc234..7e84e7466 100644 --- a/TraceLens/PerfModel/torch_op_mapping.py +++ b/TraceLens/PerfModel/torch_op_mapping.py @@ -210,7 +210,15 @@ def _categorize_torch_op_from_registry( "aiter::wrapper_fmha_v3_bwd": perf_model.aiter__fmha_v3_backward, "aiter::mha_fwd": perf_model.aiter__mha_fwd, "aiter::fmha_v3_fwd": perf_model.aiter__fmha_v3_fwd, + "aiter::fmha_v3_bwd": perf_model.aiter__fmha_v3_bwd, "aiter::mha_bwd": perf_model.aiter__mha_bwd, + # aiter varlen FlashAttention (Wan 2.x training, sglang/vLLM inference). Issue #650 / #290. + "aiter::fmha_v3_varlen_fwd": perf_model.aiter__fmha_v3_varlen_fwd, + "aiter::fmha_v3_varlen_bwd": perf_model.aiter__fmha_v3_varlen_bwd, + "aiter::wrapper_fmha_v3_varlen_fwd": perf_model.aiter__fmha_v3_varlen_forward, + "aiter::wrapper_fmha_v3_varlen_bwd": perf_model.aiter__fmha_v3_varlen_backward, + # aten dispatcher-level Flash-Attention forward (PyTorch SDPA flash backend). Issue #650. + "aten::_flash_attention_forward": perf_model.aten___flash_attention_forward, "flash_attn_3::fwd": perf_model.flash_attn_v3_forward, "vllm::unified_attention_with_output": perf_model.vllm_unified_attention_with_output, "EvoformerAttention": perf_model.evoformer_attention, diff --git a/tests/test_aiter_fmha_v3_varlen_ops.py b/tests/test_aiter_fmha_v3_varlen_ops.py new file mode 100644 index 000000000..c52aa8f0c --- /dev/null +++ b/tests/test_aiter_fmha_v3_varlen_ops.py @@ -0,0 +1,572 @@ +############################################################################### +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +# +# See LICENSE for license information. +############################################################################### + +"""Tests for aiter varlen Flash-Attention + aten flash-attention perf models. + +Closes TraceLens #650, #290, #590. Sample event payloads come from a real +Wan 2.2 T2V A14B training trace (Primus, BF16, mbs=1). +""" + +from TraceLens.PerfModel.perf_model import ( + aiter__fmha_v3_bwd, + aiter__fmha_v3_varlen_fwd, + aiter__fmha_v3_varlen_forward, + aiter__fmha_v3_varlen_bwd, + aiter__fmha_v3_varlen_backward, + aten___flash_attention_forward, +) +from TraceLens.PerfModel.extensions.attention_perf_model_extensions import ( + aiter_fmha_v3_varlen_fwd as aiter_fmha_v3_varlen_fwd_extension, +) +from TraceLens.PerfModel.torch_op_mapping import ( + categorize_torch_op, + op_to_perf_model_class_map, +) + + +def test_new_attention_ops_are_mapped(): + assert ( + op_to_perf_model_class_map["aiter::fmha_v3_varlen_fwd"] + is aiter_fmha_v3_varlen_fwd_extension + ) + assert ( + op_to_perf_model_class_map["aiter::fmha_v3_varlen_bwd"] + is aiter__fmha_v3_varlen_bwd + ) + assert ( + op_to_perf_model_class_map["aiter::wrapper_fmha_v3_varlen_fwd"] + is aiter__fmha_v3_varlen_forward + ) + assert ( + op_to_perf_model_class_map["aiter::wrapper_fmha_v3_varlen_bwd"] + is aiter__fmha_v3_varlen_backward + ) + assert op_to_perf_model_class_map["aiter::fmha_v3_bwd"] is aiter__fmha_v3_bwd + assert ( + op_to_perf_model_class_map["aten::_flash_attention_forward"] + is aten___flash_attention_forward + ) + + +def test_varlen_fwd_categorizes_as_sdpa_fwd(): + for op in ( + "aiter::fmha_v3_varlen_fwd", + "aiter::wrapper_fmha_v3_varlen_fwd", + "aten::_flash_attention_forward", + ): + row = {"name": op, "kernel_details": []} + assert categorize_torch_op(row) == "SDPA_fwd", op + + +def test_varlen_bwd_categorizes_as_sdpa_bwd(): + for op in ( + "aiter::fmha_v3_varlen_bwd", + "aiter::wrapper_fmha_v3_varlen_bwd", + "aiter::fmha_v3_bwd", + ): + row = {"name": op, "kernel_details": []} + assert categorize_torch_op(row) == "SDPA_bwd", op + + +def test_extension_mapping_falls_back_to_core_for_unannotated_varlen_fwd(): + """The extension mapping has priority, but plain training traces still use + the core shape-based SDPA model because they have no inference annotation.""" + cls = op_to_perf_model_class_map["aiter::fmha_v3_varlen_fwd"] + assert cls is aiter_fmha_v3_varlen_fwd_extension + mapped = cls(_WAN22_VARLEN_FWD) + core = aiter__fmha_v3_varlen_fwd(_WAN22_VARLEN_FWD) + assert mapped.param_details["_fallback_core"] is True + assert mapped.param_details["dtype_A_B"] == ("c10::BFloat16", "c10::BFloat16") + assert mapped.flops() == core.flops() + + +def test_extension_mapping_uses_annotation_when_available(): + event = { + **_WAN22_VARLEN_FWD, + "annotation": "attn_req_ctx_10_100_tail_a_b_c", + } + mapped = op_to_perf_model_class_map["aiter::fmha_v3_varlen_fwd"](event) + assert "_fallback_core" not in mapped.param_details + assert mapped.param_details["c_sq"] == 10 + assert mapped.param_details["c_sqsq"] == 100 + expected = 2 * 40 * (2 * 100 * 128 - 100 * 128) + assert mapped.flops() == expected + + +# Real Wan 2.2 event payloads (Primus BF16 training, mbs=1) + +_WAN22_VARLEN_FWD = { + "name": "aiter::fmha_v3_varlen_fwd", + "args": { + "Input Dims": ([[32760, 40, 128]] * 3 + [[2], [2]] + [[]] * 20), + "Input type": ["c10::BFloat16"] * 3 + ["int"] * 2 + ["Scalar"] * 13 + [""] * 7, + "Input Strides": [[5120, 128, 1]] * 3 + [[1], [1]] + [[]] * 20, + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "32760", + "32760", + "0", + "0.", + "0.088388347648318447", + "0.", + "False", + "False", + "-1", + "-1", + "True", + "False", + "1", + "", + "", + "", + "", + "", + "", + "", + ], + }, +} + +_WAN22_VARLEN_BWD = { + "name": "aiter::fmha_v3_varlen_bwd", + "args": { + "Input Dims": [ + [32760, 40, 128], + [32760, 40, 128], + [512, 40, 128], + [512, 40, 128], + [32760, 40, 128], + [40, 32760], + [2], + [2], + *([[]] * 11), + [32760, 40, 128], + [512, 40, 128], + [512, 40, 128], + [], + [2], + [], + [], + [], + ], + "Input type": ["c10::BFloat16"] * 5 + + ["float", "int", "int"] + + ["Scalar"] * 11 + + ["c10::BFloat16"] * 3 + + ["", "long int", "", "", ""], + "Input Strides": [[5120, 128, 1]] * 5 + + [[32760, 1], [1], [1]] + + [[]] * 11 + + [[5120, 128, 1]] * 3 + + [[], [1], [], [], []], + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "", + "", + "", + "32760", + "512", + "0.", + "0.088388347648318447", + "False", + "False", + "-1", + "-1", + "False", + "True", + "1", + "", + "", + "", + "", + "", + "", + "", + "", + ], + }, +} + +_WAN22_ATEN_FLASH_FWD = { + "name": "aten::_flash_attention_forward", + "args": { + "Input Dims": [ + [1, 6240, 1, 384], + [1, 6240, 1, 384], + [1, 6240, 1, 384], + [], + [], + *([[]] * 10), + ], + "Input type": ["c10::BFloat16"] * 3 + [""] * 2 + ["Scalar"] * 6 + [""] * 4, + "Input Strides": [[7188480, 1152, 7188480, 1]] * 3 + [[]] * 12, + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "6240", + "6240", + "0.", + "False", + "False", + "0.051031036307982884", + "", + "", + "", + "", + ], + }, +} + + +def test_varlen_fwd_param_extraction(): + m = aiter__fmha_v3_varlen_fwd(_WAN22_VARLEN_FWD) + p = m.param_details + assert p["B"] == 1 + assert p["N_Q"] == 32760 + assert p["N_KV"] == 32760 + assert p["H_Q"] == 40 + assert p["H_KV"] == 40 + assert p["d_h_qk"] == 128 + assert p["d_h_v"] == 128 + assert p["dropout"] == 0.0 + assert p["causal"] is False + assert p["max_seqlen_q"] == 32760.0 + assert p["max_seqlen_kv"] == 32760.0 + assert p["num_seqs_q"] == 1 + assert p["num_seqs_kv"] == 1 + assert p["dtype_A_B"] == ("c10::BFloat16", "c10::BFloat16") + + +def test_varlen_fwd_flops_matches_sdpa_formula(): + m = aiter__fmha_v3_varlen_fwd(_WAN22_VARLEN_FWD) + # Non-causal self-attention: 4 * B * H * N^2 * d_h (QK^T + PV, each 2*B*H*N^2*d) + expected = 4 * 1 * 40 * 32760**2 * 128 + assert m.flops() == expected, (m.flops(), expected) + + +def test_varlen_bwd_param_extraction_cross_attention(): + m = aiter__fmha_v3_varlen_bwd(_WAN22_VARLEN_BWD) + p = m.param_details + assert p["B"] == 1 + assert p["N_Q"] == 32760 + assert p["N_KV"] == 512 + assert p["H_Q"] == 40 + assert p["H_KV"] == 40 + assert p["d_h_qk"] == 128 + assert p["d_h_v"] == 128 + assert p["causal"] is False + assert p["max_seqlen_q"] == 32760.0 + assert p["max_seqlen_kv"] == 512.0 + assert p["dtype_A_B"] == ("c10::BFloat16", "c10::BFloat16") + + +def test_varlen_fwd_multi_seq_packing_accumulates_flops(): + """Multi-sequence packed varlen: when ``cu_seqlens`` length > 2, the FLOPs + must scale with ``num_seqs_q`` (mirror of ``flash_attention_varlen_forward``). + + Construct a 2-sequence packed varlen event: T=200 with cu_seqlens=[0,100,200], + so max_seqlen_q=100 and num_seqs_q=2. Single max-seq FLOPs: + 4 * 1 * 4 * 100^2 * 64 = 10_240_000 + Remainder seq (length (200-100)/(2-1) = 100): + 4 * 1 * 4 * 100^2 * 64 = 10_240_000 + Total = 20_480_000 (~ 2x the single-seq estimate at the same T).""" + event = { + "name": "aiter::fmha_v3_varlen_fwd", + "args": { + "Input Dims": ( + [[200, 4, 64]] * 3 # Q, K, V — packed 2 sequences of 100 + + [[3], [3]] # cu_seqlens_q, cu_seqlens_kv: 3 entries -> B=2 seqs + + [[]] * 20 + ), + "Input type": ["c10::BFloat16"] * 3 + + ["int"] * 2 + + ["Scalar"] * 13 + + [""] * 7, + "Input Strides": [[256, 64, 1]] * 3 + [[1], [1]] + [[]] * 20, + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "100", + "100", + "0", + "0.", + "0.125", + "0.", + "False", + "False", + "-1", + "-1", + "True", + "False", + "1", + "", + "", + "", + "", + "", + "", + "", + ], + }, + } + m = aiter__fmha_v3_varlen_fwd(event) + assert m.param_details["num_seqs_q"] == 2 + assert m.param_details["max_seqlen_q"] == 100.0 + expected = 4 * 1 * 4 * 100**2 * 64 + 1 * 4 * 1 * 4 * 100**2 * 64 + assert m.flops() == expected, (m.flops(), expected) + + +def test_varlen_bwd_multi_seq_packing_accumulates_flops(): + """Same multi-seq packing scenario for the backward perf model. + bwd / fwd ratio remains 5/2 for the square self-attention case.""" + bwd_event = { + "name": "aiter::fmha_v3_varlen_bwd", + "args": { + "Input Dims": ( + [[200, 4, 64]] * 5 # dout, q, k, v, out + + [[4, 200], [3], [3]] # softmax_lse, cu_seqlens_q, cu_seqlens_kv + + [[]] * 11 + + [[200, 4, 64]] * 3 # dq, dk, dv + + [[], [3], [], [], []] + ), + "Input type": ["c10::BFloat16"] * 5 + + ["float", "int", "int"] + + ["Scalar"] * 11 + + ["c10::BFloat16"] * 3 + + ["", "long int", "", "", ""], + "Input Strides": [[256, 64, 1]] * 5 + + [[200, 1], [1], [1]] + + [[]] * 11 + + [[256, 64, 1]] * 3 + + [[], [1], [], [], []], + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "", + "", + "", + "100", + "100", + "0.", + "0.125", + "False", + "False", + "-1", + "-1", + "False", + "True", + "1", + "", + "", + "", + "", + "", + "", + "", + "", + ], + }, + } + fwd_event = { + "name": "aiter::fmha_v3_varlen_fwd", + "args": { + "Input Dims": ([[200, 4, 64]] * 3 + [[3], [3]] + [[]] * 20), + "Input type": ["c10::BFloat16"] * 3 + + ["int"] * 2 + + ["Scalar"] * 13 + + [""] * 7, + "Input Strides": [[256, 64, 1]] * 3 + [[1], [1]] + [[]] * 20, + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "100", + "100", + "0", + "0.", + "0.125", + "0.", + "False", + "False", + "-1", + "-1", + "True", + "False", + "1", + "", + "", + "", + "", + "", + "", + "", + ], + }, + } + m_bwd = aiter__fmha_v3_varlen_bwd(bwd_event) + m_fwd = aiter__fmha_v3_varlen_fwd(fwd_event) + assert m_bwd.param_details["num_seqs_q"] == 2 + assert m_bwd.flops() / m_fwd.flops() == 2.5 + + +def test_varlen_bwd_fwd_flops_ratio_is_5_over_2_for_square(): + """Standard FA bwd identity: flops_bwd / flops_fwd == 5/2 for N_Q==N_KV.""" + m_fwd = aiter__fmha_v3_varlen_fwd(_WAN22_VARLEN_FWD) + sym_bwd = { + "name": "aiter::fmha_v3_varlen_bwd", + "args": { + "Input Dims": ( + [[32760, 40, 128]] * 5 + + [[40, 32760], [2], [2]] + + [[]] * 11 + + [[32760, 40, 128]] * 3 + + [[], [2], [], [], []] + ), + "Input type": ["c10::BFloat16"] * 5 + + ["float", "int", "int"] + + ["Scalar"] * 11 + + ["c10::BFloat16"] * 3 + + ["", "long int", "", "", ""], + "Input Strides": [[5120, 128, 1]] * 5 + + [[32760, 1], [1], [1]] + + [[]] * 11 + + [[5120, 128, 1]] * 3 + + [[], [1], [], [], []], + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "", + "", + "", + "32760", + "32760", + "0.", + "0.088388347648318447", + "False", + "False", + "-1", + "-1", + "False", + "True", + "1", + "", + "", + "", + "", + "", + "", + "", + "", + ], + }, + } + m_bwd = aiter__fmha_v3_varlen_bwd(sym_bwd) + assert m_bwd.flops() / m_fwd.flops() == 2.5 + + +def test_aten_flash_attention_forward_param_extraction(): + m = aten___flash_attention_forward(_WAN22_ATEN_FLASH_FWD) + p = m.param_details + assert p["B"] == 1 + assert p["N_Q"] == 6240 + assert p["N_KV"] == 6240 + assert p["H_Q"] == 1 + assert p["H_KV"] == 1 + assert p["d_h_qk"] == 384 + assert p["d_h_v"] == 384 + assert p["dropout"] == 0.0 + assert p["causal"] is False + assert p["dtype_A_B"] == ("c10::BFloat16", "c10::BFloat16") + + +def test_aten_flash_attention_forward_flops_matches_expected(): + """At the Wan 2.2 shape (B=1, S=6240, H=1, d_h=384, non-causal), + flops = 4 * B * H * N^2 * d_h = 4 * 6240^2 * 384 = 59808153600 (~59.81 GFLOPS).""" + m = aten___flash_attention_forward(_WAN22_ATEN_FLASH_FWD) + expected = 4 * 1 * 1 * 6240**2 * 384 + assert m.flops() == expected + + +def test_wrapper_varlen_fwd_indices_shift_by_one(): + """The wrapper variant adds a leading ``out`` tensor; arg indices shift by +1.""" + wrapper_event = { + "name": "aiter::wrapper_fmha_v3_varlen_fwd", + "args": { + "Input Dims": ( + [[32760, 40, 128]] + [[32760, 40, 128]] * 3 + [[2], [2]] + [[]] * 19 + ), + "Input type": ["c10::BFloat16"] * 4 + + ["int"] * 2 + + ["Scalar"] * 13 + + [""] * 6, + "Input Strides": [[5120, 128, 1]] * 4 + [[1], [1]] + [[]] * 19, + "Concrete Inputs": ( + ["", "", "", "", "", ""] + + ["32760", "32760", "0"] + + ["0.", "0.088388347648318447", "0."] + + ["False", "False", "-1", "-1", "True", "False", "1"] + + [""] * 5 + ), + }, + } + m = aiter__fmha_v3_varlen_forward(wrapper_event) + p = m.param_details + assert p["N_Q"] == 32760 and p["N_KV"] == 32760 + assert p["H_Q"] == 40 and p["d_h_qk"] == 128 + assert p["max_seqlen_q"] == 32760.0 + assert p["causal"] is False + + +def test_aiter_fmha_v3_bwd_uses_mha_bwd_layout(): + """``aiter::fmha_v3_bwd`` shares the ``aiter::mha_bwd`` argument layout.""" + qkv = [2, 512, 16, 64] + event = { + "name": "aiter::fmha_v3_bwd", + "args": { + "Input Dims": [qkv, qkv, qkv, qkv], + "Concrete Inputs": [ + "", + "", + "", + "", + "", + "", + "0.0", + "0.125", + "True", + "-1", + "-1", + "False", + ], + }, + } + m = aiter__fmha_v3_bwd(event) + p = m.param_details + assert p["B"] == 2 and p["N_Q"] == 512 and p["H_Q"] == 16 and p["d_h_qk"] == 64 + assert p["causal"] is True + assert m.flops() > 0 diff --git a/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/SDPA_fwd_intersect_mi300.csv b/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/SDPA_fwd_intersect_mi300.csv index 40840de02..4d4d00004 100644 --- a/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/SDPA_fwd_intersect_mi300.csv +++ b/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/SDPA_fwd_intersect_mi300.csv @@ -1,2 +1,3 @@ name,param: B,param: H_KV,param: H_Q,param: N_KV,param: N_Q,param: causal,param: d_h_qk,param: d_h_v,param: dropout,param: dtype_A_B,param: flash_impl,Kernel Time (µs)_sum__mi300_diff,Kernel Time (µs)_sum__mi300_pct,Kernel Time (µs)_mean__mi300_diff,Kernel Time (µs)_mean__mi300_pct,name_count__mi300_diff,name_count__mi300_pct,TFLOPS/s_mean__mi300_diff,TFLOPS/s_mean__mi300_pct,TB/s_mean__mi300_diff,TB/s_mean__mi300_pct,h100::GFLOPS_first,h100::Data Moved (MB)_first,h100::FLOPS/Byte_first,h100::TB/s_mean,h100::TB/s_median,h100::TB/s_std,h100::TB/s_min,h100::TB/s_max,h100::TFLOPS/s_mean,h100::TFLOPS/s_median,h100::TFLOPS/s_std,h100::TFLOPS/s_min,h100::TFLOPS/s_max,h100::process_name_first,h100::process_label_first,h100::thread_name_first,h100::Compute Spec,h100::kernel_details__summarize_kernel_stats,h100::trunc_kernel_details,h100::Input Dims_first,h100::Input type_first,h100::Input Strides_first,h100::Concrete Inputs_first,h100::Kernel Time (µs)_mean,h100::Kernel Time (µs)_median,h100::Kernel Time (µs)_std,h100::Kernel Time (µs)_min,h100::Kernel Time (µs)_max,h100::Kernel Time (µs)_sum,h100::name_count,h100::UID_first,mi300::TB/s_mean,mi300::TB/s_median,mi300::TB/s_std,mi300::TB/s_min,mi300::TB/s_max,mi300::TFLOPS/s_mean,mi300::TFLOPS/s_median,mi300::TFLOPS/s_std,mi300::TFLOPS/s_min,mi300::TFLOPS/s_max,mi300::process_name_first,mi300::process_label_first,mi300::thread_name_first,mi300::Compute Spec,mi300::kernel_details__summarize_kernel_stats,mi300::trunc_kernel_details,mi300::Kernel Time (µs)_mean,mi300::Kernel Time (µs)_median,mi300::Kernel Time (µs)_std,mi300::Kernel Time (µs)_min,mi300::Kernel Time (µs)_max,mi300::Kernel Time (µs)_sum,mi300::name_count,mi300::UID_first +aten::_flash_attention_forward,16,16,16,559,559,True,64,64,0.0,,True,1293.638916015625,47.51301470015281,53.90162150065103,47.513014700152794,0,0.0,-28.85641375823659,-31.95997550348616,-0.20648596606967173,-31.959975503486188,10.239377408,69.875,139.75,0.6460767344679231,0.6439839021131257,0.0122887818060238,0.6123786199798398,0.6704724410436742,90.28922364189224,89.9967503203093,1.7173572573918277,85.57991214218261,93.69852363585343,python3,CPU,thread 25286 (python3),,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel, std::array >(int, at::native::FillFunctor, std::array)', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(107.545), 'mean_duration_us': np.float64(4.481041666666667), 'median_duration_us': np.float64(4.69), 'std_dev_duration_us': np.float64(0.9520142890719772), 'min_duration_us': np.float64(0.681), 'max_duration_us': np.float64(5.411)}, {'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Fi...', 'stream': 0, 'mean_duration_us': np.float64(4.48)}, {'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]",167.3476359049479,168.906494140625,10.235203983555548,139.55908203125,177.92724609375,4016.34326171875,24,165 aten::_scaled_dot_product_flash_attention,16,16,16,559,559,True,64,64,0.0,"('c10::BFloat16', 'c10::BFloat16')",True,1293.638916015625,47.51301470015281,53.90162150065103,47.513014700152794,0,0.0,-28.85641375823659,-31.95997550348616,-0.20648596606967173,-31.959975503486188,10.239377408,69.875,139.75,0.6460767344679231,0.6439839021131257,0.0122887818060238,0.6123786199798398,0.6704724410436742,90.28922364189224,89.9967503203093,1.7173572573918277,85.57991214218261,93.69852363585343,python3,CPU,thread 25286 (python3),matrix_bf16,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel, std::array >(int, at::native::FillFunctor, std::array)', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(107.545), 'mean_duration_us': np.float64(4.481041666666667), 'median_duration_us': np.float64(4.69), 'std_dev_duration_us': np.float64(0.9520142890719772), 'min_duration_us': np.float64(0.681), 'max_duration_us': np.float64(5.411)}, {'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Fi...', 'stream': 0, 'mean_duration_us': np.float64(4.48)}, {'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]",167.3476359049479,168.906494140625,10.235203983555548,139.55908203125,177.92724609375,4016.34326171875,24,158 diff --git a/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_summary.csv b/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_summary.csv index 4869aae2a..4285cfcc2 100644 --- a/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_summary.csv +++ b/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_summary.csv @@ -2,7 +2,7 @@ name,total_direct_kernel_time_ms__mi300_diff,total_direct_kernel_time_ms__mi300_ aten::mul,-2.2116113281250005,-32.75054213905065,0.0,0.0,6752.89990234375,6.75289990234375,6.75289990234375,220.0,['elementwise'],21.165327193718745,21.165327193718745,4541.28857421875,4.54128857421875,4.54128857421875,220,['elementwise'],15.267951881722706 aten::mm,1.99493359375,29.914962225772584,0.0,0.0,6668.681640625,6.668681640625,6.668681640625,96.0,['GEMM'],20.901365474939983,42.06669266865873,8663.615234375,8.663615234375,8.663615234375,96,['GEMM'],29.12734092062252 aten::copy_,-1.9112404785156252,-37.461091684302325,0.0,0.0,5101.934814453125,5.101934814453125,5.101934814453125,197.0,['elementwise'],15.99077747790201,58.057470146560746,3190.6943359375,3.1906943359375,3.1906943359375,197,['elementwise'],10.727212506806964 -aten::_flash_attention_forward,1.186093994140625,43.56308447564189,0.0,0.0,2722.704345703125,2.722704345703125,2.722704345703125,24.0,['other'],8.53365652711153,66.59112667367228,4016.34326171875,3.90879833984375,4.01634326171875,24,['other'],13.141500257635307 +aten::_flash_attention_forward,1.186093994140625,43.56308447564189,0.0,0.0,2722.704345703125,2.722704345703125,2.722704345703125,24.0,['SDPA_fwd'],8.53365652711153,66.59112667367228,4016.34326171875,3.90879833984375,4.01634326171875,24,['SDPA_fwd'],13.141500257635307 aten::add,-0.71191162109375,-27.302300610596138,0.0,0.0,2607.51513671875,2.60751513671875,2.60751513671875,145.0,['elementwise'],8.17262388445474,74.76375055812701,1895.603515625,1.895603515625,1.895603515625,145,['elementwise'],6.373077330450393 aten::addmm,0.625686767578125,26.864163599284897,0.0,0.0,2329.075927734375,2.329075927734375,2.329075927734375,72.0,['GEMM'],7.29992370424488,82.06367426237189,2954.7626953125,2.9547626953125,2.9547626953125,72,['GEMM'],9.934003073500232 aten::cat,-0.6617131347656249,-32.7940689467756,0.0,0.0,2017.782958984375,2.017782958984375,2.017782958984375,49.0,['other'],6.324251380949954,88.38792564332185,1356.06982421875,1.35606982421875,1.35606982421875,49,['other'],4.559148463272868 diff --git a/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_unique_args_intersect_mi300.csv b/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_unique_args_intersect_mi300.csv index 3c68bbdb4..38b59f7a2 100644 --- a/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_unique_args_intersect_mi300.csv +++ b/tests/traces/compare_test_e2e/reference/compare_h100_mi300_qwen_csvs/ops_unique_args_intersect_mi300.csv @@ -1,6 +1,6 @@ name,Input type,Input Dims,Input Strides,Concrete Inputs,total_direct_kernel_time_sum__mi300_diff,total_direct_kernel_time_sum__mi300_pct,total_direct_kernel_time_mean__mi300_diff,total_direct_kernel_time_mean__mi300_pct,operation_count__mi300_diff,operation_count__mi300_pct,h100::op category,h100::process_name,h100::process_label,h100::thread_name,h100::operation_count,h100::total_direct_kernel_time_mean,h100::total_subtree_kernel_time_mean,h100::total_direct_kernel_time_median,h100::total_subtree_kernel_time_median,h100::total_direct_kernel_time_std,h100::total_subtree_kernel_time_std,h100::total_direct_kernel_time_min,h100::total_subtree_kernel_time_min,h100::total_direct_kernel_time_max,h100::total_subtree_kernel_time_max,h100::total_direct_kernel_time_sum,h100::total_subtree_kernel_time_sum,h100::ex_UID,h100::kernel_details_summary,h100::trunc_kernel_details,h100::Percentage (%),h100::Cumulative Percentage (%),mi300::op category,mi300::process_name,mi300::process_label,mi300::thread_name,mi300::operation_count,mi300::total_direct_kernel_time_mean,mi300::total_subtree_kernel_time_mean,mi300::total_direct_kernel_time_median,mi300::total_subtree_kernel_time_median,mi300::total_direct_kernel_time_std,mi300::total_subtree_kernel_time_std,mi300::total_direct_kernel_time_min,mi300::total_subtree_kernel_time_min,mi300::total_direct_kernel_time_max,mi300::total_subtree_kernel_time_max,mi300::total_direct_kernel_time_sum,mi300::total_subtree_kernel_time_sum,mi300::ex_UID,mi300::kernel_details_summary,mi300::trunc_kernel_details,mi300::Percentage (%),mi300::Cumulative Percentage (%) aten::mm,"('c10::BFloat16', 'c10::BFloat16')","((8944, 1024), (1024, 2816))","((1024, 1), (1, 1024))","('', '')",1253.02099609375,30.898717507688296,26.104604085286482,30.89871750768832,0.0,0.0,GEMM,python3,CPU,thread 25286 (python3),48.0,84.4844258626302,84.4844258626302,84.4635009765625,84.4635009765625,1.1997237297511991,1.1997237297511991,82.27197265625,82.27197265625,86.912109375,86.912109375,4055.25244140625,4055.25244140625,215.0,"[{'name': 'Memset (Device)', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(39.75), 'mean_duration_us': np.float64(0.828125), 'median_duration_us': np.float64(0.8), 'std_dev_duration_us': np.float64(0.07937112851450877), 'min_duration_us': np.float64(0.768), 'max_duration_us': np.float64(1.057)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(4015.5009999999997), 'mean_duration_us': np.float64(83.65627083333332), 'median_duration_us': np.float64(83.679), 'std_dev_duration_us': np.float64(1.1791844593401435), 'min_duration_us': np.float64(81.472), 'max_duration_us': np.float64(86.016)}]","[{'name': 'Memset (Device)', 'stream': 7, 'mean_duration_us': np.float64(0.83)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warp...', 'stream': 7, 'mean_duration_us': np.float64(83.66)}]",12.710205395714588,12.710205395714588,GEMM,python3,CPU,thread 23232 (python3),48,110.58902994791669,110.58902994791669,110.452880859375,110.452880859375,3.280588398059727,3.280588398059727,105.0,105.0,118.55078125,118.55078125,5308.2734375,5308.2734375,231,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_7_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(5308.273), 'mean_duration_us': np.float64(110.58902083333334), 'median_duration_us': np.float64(110.453), 'std_dev_duration_us': np.float64(3.2462516941953186), 'min_duration_us': np.float64(105.0), 'max_duration_us': np.float64(118.551)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(110.59)}]",17.846578585400604,17.846578585400604 -aten::_flash_attention_forward,"('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",1186.093994140625,43.56308447564189,49.42058308919272,43.563084475641894,0.0,0.0,other,python3,CPU,thread 25286 (python3),24.0,113.44601440429688,113.44601440429688,113.7750244140625,113.7750244140625,2.173453067400901,2.173453067400901,109.280029296875,109.280029296875,119.64697265625,119.64697265625,2722.704345703125,2722.704345703125,164.0,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 72, 'total_duration_us': np.float64(2108.187), 'mean_duration_us': np.float64(29.280375), 'median_duration_us': np.float64(29.312), 'std_dev_duration_us': np.float64(0.3210814381594732), 'min_duration_us': np.float64(28.672), 'max_duration_us': np.float64(30.048)}]","[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(29.28)}]",6.607603257237321,35.151388884308325,elementwise,python3,CPU,thread 23232 (python3),72,15.11367458767361,15.11367458767361,15.19287109375,15.19287109375,0.5665176534834554,0.5665176534834554,12.98779296875,12.98779296875,16.27587890625,16.27587890625,1088.1845703125,1088.1845703125,146,"[{'name': 'void at::native::elementwise_kernel_manual_unroll<128, 8, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int, bool)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int, bool)#1})', 'stream': 0, 'count': 72, 'total_duration_us': np.float64(1088.184), 'mean_duration_us': np.float64(15.113666666666667), 'median_duration_us': np.float64(15.193), 'std_dev_duration_us': np.float64(0.5625623422242988), 'min_duration_us': np.float64(12.988), 'max_duration_us': np.float64(16.276)}]","[{'name': 'void at::native::elementwise_kernel_manual_unroll<128, 8, at::na...', 'stream': 0, 'mean_duration_us': np.float64(15.11)}]",3.658509999185101,60.98291891379351 aten::cat,"('TensorList', 'Scalar')","(((16, 16, 559, 32), (16, 16, 559, 32)), ())","(((286208, 32, 512, 1), (572416, 64, 1024, 1)), ())","('', '-1')",-664.590087890625,-32.99315947540343,-13.845626831054684,-32.993159475403424,0.0,0.0,other,python3,CPU,thread 25286 (python3),48.0,41.96514383951823,41.96514383951823,41.9678955078125,41.9678955078125,0.2933502450882166,0.2933502450882166,41.3759765625,41.3759765625,42.464111328125,42.464111328125,2014.326904296875,2014.326904296875,129.0,"[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(2014.326), 'mean_duration_us': np.float64(41.965125), 'median_duration_us': np.float64(41.968), 'std_dev_duration_us': np.float64(0.290296301575362), 'min_duration_us': np.float64(41.376), 'max_duration_us': np.float64(42.464)}]","[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(1349.738), 'mean_duration_us': np.float64(28.119541666666667), 'median_duration_us': np.float64(27.903), 'std_dev_duration_us': np.float64(0.8283528575415325), 'min_duration_us': np.float64(26.82), 'max_duration_us': np.float64(30.348)}]","[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel > >(at::TensorIteratorBase&, at::native::BinaryFunctor > const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast > >(at::TensorIteratorBase&, at::native::BinaryFunctor > const&)::{lambda(int)#1})",7,3913.711,145,26.991110344827586,23.392,32.672,8.614304613826809,11.216508107775581 elementwise,aten::copy_,"void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})",7,2766.267,96,28.815281249999998,26.913,30.048,6.088713903805581,7.927988610751288 -other,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,2722.705,24,113.44604166666666,109.28,119.647,5.992831418464297,7.803142007057015 +SDPA_fwd,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,2722.705,24,113.44604166666666,109.28,119.647,5.992831418464297,7.803142007057015 GEMM,aten::addmm,sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas,7,2266.292,72,31.476277777777778,30.912,32.256,4.988239967610992,6.4950842289037025 other,aten::cat,"void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)",7,2014.326,48,41.965125,41.376,42.464,4.433648206408521,5.772961751826632 elementwise,aten::add,"void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast >(at::TensorIteratorBase&, at::native::CUDAFunctor_add const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast >(at::TensorIteratorBase&, at::native::CUDAFunctor_add const&)::{lambda(int)#1})",7,1671.096,48,34.8145,34.271,35.615,3.6781790947127995,4.789281025827238 diff --git a/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary.csv b/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary.csv index 1c0587fec..ea93a42fc 100644 --- a/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary.csv +++ b/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary.csv @@ -2,7 +2,7 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_ker aten::mul,6752.89990234375,6752.89990234375,6.75289990234375,6.75289990234375,220,['elementwise'],21.165327193718745,21.165327193718745 aten::mm,6668.681640625,6668.681640625,6.668681640625,6.668681640625,96,['GEMM'],20.901365474939983,42.06669266865873 aten::copy_,5101.934814453125,5101.934814453125,5.101934814453125,5.101934814453125,197,['elementwise'],15.99077747790201,58.057470146560746 -aten::_flash_attention_forward,2722.704345703125,2722.704345703125,2.722704345703125,2.722704345703125,24,['other'],8.53365652711153,66.59112667367228 +aten::_flash_attention_forward,2722.704345703125,2722.704345703125,2.722704345703125,2.722704345703125,24,['SDPA_fwd'],8.53365652711153,66.59112667367228 aten::add,2607.51513671875,2607.51513671875,2.60751513671875,2.60751513671875,145,['elementwise'],8.17262388445474,74.76375055812701 aten::addmm,2329.075927734375,2329.075927734375,2.329075927734375,2.329075927734375,72,['GEMM'],7.29992370424488,82.06367426237189 aten::cat,2017.782958984375,2017.782958984375,2.017782958984375,2.017782958984375,49,['other'],6.324251380949954,88.38792564332185 diff --git a/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary_by_category.csv index 8a3d282f4..1e19d5b37 100644 --- a/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_summary_by_category.csv @@ -1,5 +1,6 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) elementwise,735,17.360615234375,54.41263857506248,54.41263857506248 GEMM,169,9.001277587890625,28.212321826815455,82.62496040187793 -other,76,4.796135009765625,15.032322145291522,97.65728254716944 -reduce,50,0.747455322265625,2.3427174528305534,100.0 +SDPA_fwd,24,2.722704345703125,8.53365652711153,91.15861692898946 +other,52,2.0734306640625,6.498665618179992,97.65728254716946 +reduce,50,0.747455322265625,2.3427174528305534,100.00000000000001 diff --git a/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_unique_args.csv b/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_unique_args.csv index 5cc0743f5..cc31a0fae 100644 --- a/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/compare_test_e2e/reference/h100_perf_report_csvs/ops_unique_args.csv @@ -1,6 +1,6 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::mm,GEMM,python3,CPU,thread 25286 (python3),"((8944, 1024), (1024, 2816))","('c10::BFloat16', 'c10::BFloat16')","((1024, 1), (1, 1024))","('', '')",48,84.4844258626302,84.4844258626302,84.4635009765625,84.4635009765625,1.1997237297511993,1.1997237297511993,82.27197265625,82.27197265625,86.912109375,86.912109375,4055.25244140625,4055.25244140625,215,"[{'name': 'Memset (Device)', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(39.75), 'mean_duration_us': np.float64(0.828125), 'median_duration_us': np.float64(0.8), 'std_dev_duration_us': np.float64(0.07937112851450877), 'min_duration_us': np.float64(0.768), 'max_duration_us': np.float64(1.057)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(4015.5009999999997), 'mean_duration_us': np.float64(83.65627083333332), 'median_duration_us': np.float64(83.679), 'std_dev_duration_us': np.float64(1.1791844593401435), 'min_duration_us': np.float64(81.472), 'max_duration_us': np.float64(86.016)}]","[{'name': 'Memset (Device)', 'stream': 7, 'mean_duration_us': np.float64(0.83)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warp...', 'stream': 7, 'mean_duration_us': np.float64(83.66)}]",12.710205395714588,12.710205395714588 -aten::_flash_attention_forward,other,python3,CPU,thread 25286 (python3),"((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",24,113.44601440429688,113.44601440429688,113.7750244140625,113.7750244140625,2.173453067400901,2.173453067400901,109.280029296875,109.280029296875,119.64697265625,119.64697265625,2722.704345703125,2722.704345703125,164,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 72, 'total_duration_us': np.float64(2108.187), 'mean_duration_us': np.float64(29.280375), 'median_duration_us': np.float64(29.312), 'std_dev_duration_us': np.float64(0.3210814381594732), 'min_duration_us': np.float64(28.672), 'max_duration_us': np.float64(30.048)}]","[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(29.28)}]",6.607603257237321,35.151388884308325 aten::cat,other,python3,CPU,thread 25286 (python3),"(((16, 16, 559, 32), (16, 16, 559, 32)), ())","('TensorList', 'Scalar')","(((286208, 32, 512, 1), (572416, 64, 1024, 1)), ())","('', '-1')",48,41.96514383951823,41.96514383951823,41.9678955078125,41.9678955078125,0.2933502450882166,0.2933502450882166,41.3759765625,41.3759765625,42.464111328125,42.464111328125,2014.326904296875,2014.326904296875,129,"[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(2014.326), 'mean_duration_us': np.float64(41.965125), 'median_duration_us': np.float64(41.968), 'std_dev_duration_us': np.float64(0.290296301575362), 'min_duration_us': np.float64(41.376), 'max_duration_us': np.float64(42.464)}]","[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, std::array >(int, at::native::FillFunctor, std::array)', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(107.545), 'mean_duration_us': np.float64(4.481041666666667), 'median_duration_us': np.float64(4.69), 'std_dev_duration_us': np.float64(0.9520142890719772), 'min_duration_us': np.float64(0.681), 'max_duration_us': np.float64(5.411)}, {'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Fi...', 'stream': 0, 'mean_duration_us': np.float64(4.48)}, {'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]","[[16, 16, 559, 64], [16, 16, 559, 64], [16, 16, 559, 64], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar', 'Scalar', 'Scalar']","[[572416, 35776, 64, 1], [572416, 35776, 64, 1], [572416, 35776, 64, 1], [], [], [], []]","['', '', '', '0.', 'True', 'False', '0.125']",167.3476359049479,168.906494140625,10.235203983555547,139.55908203125,177.92724609375,4016.34326171875,24,158 +aten::_flash_attention_forward,16,559,16,559,16,64,64,0.0,True,True,,10.239377408,69.875,139.75,0.43959076839825145,0.43379994577448006,0.03016579398980997,0.4117933009618683,0.525005230283713,61.432809883655636,60.62354242198358,4.215669710075947,57.54811380942109,73.36948093214889,python3,CPU,thread 23232 (python3),,"[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor, std::array >(int, at::native::FillFunctor, std::array)', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(107.545), 'mean_duration_us': np.float64(4.481041666666667), 'median_duration_us': np.float64(4.69), 'std_dev_duration_us': np.float64(0.9520142890719772), 'min_duration_us': np.float64(0.681), 'max_duration_us': np.float64(5.411)}, {'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Fi...', 'stream': 0, 'mean_duration_us': np.float64(4.48)}, {'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]","[[16, 559, 16, 64], [16, 559, 16, 64], [16, 559, 16, 64], [], [], [], [], [], [], [], [], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '']","[[572416, 64, 35776, 1], [572416, 64, 35776, 1], [572416, 64, 35776, 1], [], [], [], [], [], [], [], [], [], [], [], []]","['', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '']",167.3476359049479,168.906494140625,10.235203983555547,139.55908203125,177.92724609375,4016.34326171875,24,165 diff --git a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/kernel_summary.csv b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/kernel_summary.csv index b0fd44a33..6d952de21 100644 --- a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/kernel_summary.csv +++ b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/kernel_summary.csv @@ -1,6 +1,6 @@ Parent op category,Parent cpu_op,Kernel name,Kernel stream,Kernel duration (µs)_sum,Kernel duration (µs)_count,Kernel duration (µs)_mean,Kernel duration (µs)_min,Kernel duration (µs)_max,Percent of kernels time (%),Percent of total time (%) GEMM,aten::mm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_7_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0,5308.273,48,110.58902083333334,105.0,118.551,12.327035025417512,16.021911706710675 -other,aten::_flash_attention_forward,attn_fwd,0,3908.798,24,162.86658333333332,138.557,173.077,9.07713108449432,11.797889150646034 +SDPA_fwd,aten::_flash_attention_forward,attn_fwd,0,3908.798,24,162.86658333333332,138.557,173.077,9.07713108449432,11.797889150646034 GEMM,aten::mm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0,3355.344,48,69.903,39.209,105.883,7.791883162438047,10.127404018904345 GEMM,aten::addmm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0,2954.763,72,41.038375,39.57,45.023,6.861641628606464,8.918334060862271 ,,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0,2787.634,52,53.608346153846156,39.811,102.314,6.473529518177516,8.41389013312328 diff --git a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary.csv b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary.csv index 213411270..0c046937b 100644 --- a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary.csv +++ b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary.csv @@ -1,7 +1,7 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::mm,8663.615234375,8663.615234375,8.663615234375,8.663615234375,96,['GEMM'],29.12734092062252,29.12734092062252 aten::mul,4541.28857421875,4541.28857421875,4.54128857421875,4.54128857421875,220,['elementwise'],15.267951881722706,44.395292802345224 -aten::_flash_attention_forward,3908.79833984375,4016.34326171875,3.90879833984375,4.01634326171875,24,['other'],13.141500257635307,57.53679305998053 +aten::_flash_attention_forward,3908.79833984375,4016.34326171875,3.90879833984375,4.01634326171875,24,['SDPA_fwd'],13.141500257635307,57.53679305998053 aten::copy_,3190.6943359375,3190.6943359375,3.1906943359375,3.1906943359375,197,['elementwise'],10.727212506806964,68.26400556678749 aten::addmm,2954.7626953125,2954.7626953125,2.9547626953125,2.9547626953125,72,['GEMM'],9.934003073500232,78.19800864028772 aten::add,1895.603515625,1895.603515625,1.895603515625,1.895603515625,145,['elementwise'],6.373077330450393,84.57108597073812 diff --git a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary_by_category.csv index b0c6b47d4..2cd86061c 100644 --- a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_summary_by_category.csv @@ -1,5 +1,6 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) elementwise,759,12.18791064453125,40.97613050090121,40.97613050090121 GEMM,169,11.6457197265625,39.153267956281255,80.12939845718248 -other,76,5.288919921875,17.78151147053321,97.9109099277157 -reduce,50,0.62137744140625,2.0890900722843178,100.00000000000001 +SDPA_fwd,24,3.90879833984375,13.141500257635307,93.27089871481778 +other,52,1.38012158203125,4.640011212897906,97.91090992771568 +reduce,50,0.62137744140625,2.0890900722843178,100.0 diff --git a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_unique_args.csv b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_unique_args.csv index 5fc72f823..df0d0b388 100644 --- a/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/compare_test_e2e/reference/mi300_perf_report_csvs/ops_unique_args.csv @@ -1,6 +1,6 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::mm,GEMM,python3,CPU,thread 23232 (python3),"((8944, 1024), (1024, 2816))","('c10::BFloat16', 'c10::BFloat16')","((1024, 1), (1, 1024))","('', '')",48,110.58902994791667,110.58902994791667,110.452880859375,110.452880859375,3.280588398059727,3.280588398059727,105.0,105.0,118.55078125,118.55078125,5308.2734375,5308.2734375,231,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_7_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(5308.273), 'mean_duration_us': np.float64(110.58902083333334), 'median_duration_us': np.float64(110.453), 'std_dev_duration_us': np.float64(3.2462516941953186), 'min_duration_us': np.float64(105.0), 'max_duration_us': np.float64(118.551)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(110.59)}]",17.846578585400604,17.846578585400604 -aten::_flash_attention_forward,other,python3,CPU,thread 23232 (python3),"((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",24,162.8665974934896,167.3476359049479,164.237060546875,168.906494140625,9.452351716739908,10.235203983555547,138.55712890625,139.55908203125,173.0771484375,177.92724609375,3908.79833984375,4016.34326171875,165,"[{'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]",13.141500257635306,30.988078843035908 +aten::_flash_attention_forward,SDPA_fwd,python3,CPU,thread 23232 (python3),"((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",24,162.8665974934896,167.3476359049479,164.237060546875,168.906494140625,9.452351716739908,10.235203983555547,138.55712890625,139.55908203125,173.0771484375,177.92724609375,3908.79833984375,4016.34326171875,165,"[{'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]",13.141500257635306,30.988078843035908 aten::addmm,GEMM,python3,CPU,thread 23232 (python3),"((1024,), (8944, 1024), (1024, 1024), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (1024, 1), (1, 1024), (), ())","('', '', '', '1', '1')",72,41.038370768229164,41.038370768229164,40.71240234375,40.71240234375,1.1213617423390887,1.1213617423390887,39.56982421875,39.56982421875,45.02294921875,45.02294921875,2954.7626953125,2954.7626953125,90,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 72, 'total_duration_us': np.float64(2954.763), 'mean_duration_us': np.float64(41.038375), 'median_duration_us': np.float64(40.7125), 'std_dev_duration_us': np.float64(1.1135532272552386), 'min_duration_us': np.float64(39.57), 'max_duration_us': np.float64(45.023)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(41.04)}]",9.93400307350023,40.922081916536136 aten::mm,GEMM,python3,CPU,thread 23232 (python3),"((8944, 2816), (2816, 1024))","('c10::BFloat16', 'c10::BFloat16')","((2816, 1), (1, 2816))","('', '')",24,98.94974772135417,98.94974772135417,98.324951171875,98.324951171875,3.3849417047301227,3.3849417047301227,93.65380859375,93.65380859375,105.8828125,105.8828125,2374.7939453125,2374.7939453125,253,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(2374.795), 'mean_duration_us': np.float64(98.94979166666667), 'median_duration_us': np.float64(98.32499999999999), 'std_dev_duration_us': np.float64(3.3136632546066833), 'min_duration_us': np.float64(93.654), 'max_duration_us': np.float64(105.883)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(98.95)}]",7.9841302955021485,48.906212212038284 aten::cat,other,python3,CPU,thread 23232 (python3),"(((16, 16, 559, 32), (16, 16, 559, 32)), ())","('TensorList', 'Scalar')","(((286208, 32, 512, 1), (572416, 64, 1024, 1)), ())","('', '-1')",48,28.119517008463543,28.119517008463543,27.90283203125,27.90283203125,0.8371311008823814,0.8371311008823814,26.81982421875,26.81982421875,30.34814453125,30.34814453125,1349.73681640625,1349.73681640625,130,"[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(1349.738), 'mean_duration_us': np.float64(28.119541666666667), 'median_duration_us': np.float64(27.903), 'std_dev_duration_us': np.float64(0.8283528575415325), 'min_duration_us': np.float64(26.82), 'max_duration_us': np.float64(30.348)}]","[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(122.848), 'mean_duration_us': np.float64(10.237333333333334), 'median_duration_us': np.float64(10.224), 'std_dev_duration_us': np.float64(0.21878045819699932), 'min_duration_us': np.float64(10.016), 'max_duration_us': np.float64(10.848)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(122.848), 'mean_duration_us': np.float64(10.237333333333334), 'median_duration_us': np.float64(10.224), 'std_dev_duration_us': np.float64(0.21878045819699932), 'min_duration_us': np.float64(10.016), 'max_duration_us': np.float64(10.848)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_128x64_64x6_tn_align8::Params),7,399.136,48,8.315333333333333,8.128,8.64,32.691252377293964,5.279685669563507 GEMM,aten::addmm,sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas,7,166.59,12,13.8825,13.76,14.176,13.644561586861121,2.203616901739218 GEMM,aten::addmm,sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x64x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas,7,162.432,12,13.536,13.344,13.952,13.304000406249028,2.148615766752534 -other,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,122.848,12,10.237333333333334,10.016,10.848,10.06187107162924,1.6250070781250945 +SDPA_fwd,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,122.848,12,10.237333333333334,10.016,10.848,10.06187107162924,1.6250070781250945 NORM_fwd,aten::native_layer_norm,"void at::native::(anonymous namespace)::vectorized_layer_norm_kernel(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)",7,111.20100000000001,25,4.448040000000001,4.256,4.64,9.107923002704506,1.470943052345896 CONV_fwd,aten::cudnn_convolution,void cutlass__5x_cudnn::Kernel(cutlass_tensorop_bf16_s16816fprop_optimized_bf16_256x64_32x4_nhwc_align8::Params),7,98.176,1,98.176,98.176,98.176,8.041109780609144,1.2986511371940062 elementwise,aten::gelu,"void at::native::vectorized_elementwise_kernel<4, at::native::GeluCUDAKernelImpl(at::TensorIteratorBase&, at::native::GeluType)::{lambda()#2}::operator()() const::{lambda()#4}::operator()() const::{lambda(c10::BFloat16)#1}, std::array >(int, at::native::GeluCUDAKernelImpl(at::TensorIteratorBase&, at::native::GeluType)::{lambda()#2}::operator()() const::{lambda()#4}::operator()() const::{lambda(c10::BFloat16)#1}, std::array)",7,60.542,12,5.045166666666667,4.991,5.12,4.958695285381752,0.8008366316411294 diff --git a/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv b/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv index d63331fd4..a0f46b586 100644 --- a/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv +++ b/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv @@ -1,7 +1,7 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::addmm,735.03662109375,735.03662109375,0.73503662109375,0.73503662109375,73,['GEMM'],59.98786578240299,59.98786578240299 aten::cudnn_convolution,125.632080078125,125.632080078125,0.125632080078125,0.125632080078125,1,['CONV_fwd'],10.253095072292243,70.24096085469523 -aten::_flash_attention_forward,122.84814453125,122.84814453125,0.12284814453125,0.12284814453125,12,['other'],10.025892308320707,80.26685316301594 +aten::_flash_attention_forward,122.84814453125,122.84814453125,0.12284814453125,0.12284814453125,12,['SDPA_fwd'],10.025892308320707,80.26685316301594 aten::native_layer_norm,111.200927734375,111.200927734375,0.111200927734375,0.111200927734375,25,['NORM_fwd'],9.075338746907917,89.34219190992386 aten::gelu,60.541259765625,60.541259765625,0.060541259765625,0.060541259765625,12,['elementwise'],4.94089799187665,94.28308990180051 aten::add,57.31298828125,57.31298828125,0.05731298828125,0.05731298828125,25,['elementwise'],4.677432048879577,98.96052195068009 diff --git a/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv index b55b6c90c..16049ec8f 100644 --- a/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv @@ -1,6 +1,7 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) GEMM,73,0.73503662109375,59.98786578240299,59.98786578240299 -other,13,0.130112060546875,10.61871558609367,70.60658136849666 -CONV_fwd,1,0.125632080078125,10.253095072292243,80.8596764407889 -elementwise,39,0.1233271484375,10.06498481230318,90.92466125309208 -NORM_fwd,25,0.111200927734375,9.075338746907917,100.0 +CONV_fwd,1,0.125632080078125,10.253095072292243,70.24096085469523 +elementwise,39,0.1233271484375,10.06498481230318,80.30594566699841 +SDPA_fwd,12,0.12284814453125,10.025892308320707,90.33183797531912 +NORM_fwd,25,0.111200927734375,9.075338746907917,99.40717672222704 +other,1,0.007263916015625,0.5928232777729626,100.0 diff --git a/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv b/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv index cc695523a..f1adf00eb 100644 --- a/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/h100/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv @@ -3,7 +3,7 @@ aten::addmm,GEMM,python3,CPU,thread 5617 (python3),"((768,), (788, 768), (768, 7 aten::addmm,GEMM,python3,CPU,thread 5617 (python3),"((3072,), (788, 768), (768, 3072), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (768, 1), (1, 768), (), ())","('', '', '', '1', '1')",12,13.882466634114584,13.882466634114584,13.887451171875,13.887451171875,0.1063651301764554,0.1063651301764554,13.760009765625,13.760009765625,14.176025390625,14.176025390625,166.589599609375,166.589599609375,104,"[{'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(166.59), 'mean_duration_us': np.float64(13.8825), 'median_duration_us': np.float64(13.8875), 'std_dev_duration_us': np.float64(0.10183278777813505), 'min_duration_us': np.float64(13.76), 'max_duration_us': np.float64(14.176)}]","[{'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warp...', 'stream': 7, 'mean_duration_us': np.float64(13.88)}]",13.595723335853824,46.1701002119005 aten::addmm,GEMM,python3,CPU,thread 5617 (python3),"((768,), (788, 3072), (3072, 768), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (3072, 1), (1, 3072), (), ())","('', '', '', '1', '1')",12,13.53594970703125,13.53594970703125,13.4879150390625,13.4879150390625,0.17737765723261711,0.17737765723261711,13.343994140625,13.343994140625,13.951904296875,13.951904296875,162.431396484375,162.431396484375,114,"[{'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x64x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(162.43200000000002), 'mean_duration_us': np.float64(13.536000000000001), 'median_duration_us': np.float64(13.488), 'std_dev_duration_us': np.float64(0.16983128883296716), 'min_duration_us': np.float64(13.344), 'max_duration_us': np.float64(13.952)}]","[{'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x64x64_warpg...', 'stream': 7, 'mean_duration_us': np.float64(13.54)}]",13.25636373961045,59.42646395151095 aten::cudnn_convolution,CONV_fwd,python3,CPU,thread 5617 (python3),"((4, 3, 224, 224), (768, 3, 16, 16), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'ScalarList', 'ScalarList', 'ScalarList', 'Scalar', 'Scalar', 'Scalar', 'Scalar')","((150528, 50176, 224, 1), (768, 256, 16, 1), (), (), (), (), (), (), ())","('', '', '[0, 0]', '[16, 16]', '[1, 1]', '1', 'False', 'False', 'True')",1,125.632080078125,125.632080078125,125.632080078125,125.632080078125,,,125.632080078125,125.632080078125,125.632080078125,125.632080078125,125.632080078125,125.632080078125,4,"[{'name': 'Memset (Device)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(2.112), 'mean_duration_us': np.float64(2.112), 'median_duration_us': np.float64(2.112), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(2.112), 'max_duration_us': np.float64(2.112)}, {'name': 'void cudnn::engines_precompiled::nhwcToNchwKernel<__nv_bfloat16, __nv_bfloat16, float, true, false, (cudnnKernelDataType_t)0>(cudnn::engines_precompiled::nhwc2nchw_params_t, __nv_bfloat16 const*, __nv_bfloat16*)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(2.208), 'mean_duration_us': np.float64(2.208), 'median_duration_us': np.float64(2.208), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(2.208), 'max_duration_us': np.float64(2.208)}, {'name': 'Memset (Device)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(2.272), 'mean_duration_us': np.float64(2.272), 'median_duration_us': np.float64(2.272), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(2.272), 'max_duration_us': np.float64(2.272)}, {'name': 'void cudnn::engines_precompiled::nchwToNhwcKernel<__nv_bfloat16, __nv_bfloat16, float, false, true, (cudnnKernelDataType_t)0>(cudnn::engines_precompiled::nchw2nhwc_params_t, __nv_bfloat16 const*, __nv_bfloat16*)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(10.4), 'mean_duration_us': np.float64(10.4), 'median_duration_us': np.float64(10.4), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(10.4), 'max_duration_us': np.float64(10.4)}, {'name': 'void cudnn::engines_precompiled::nchwToNhwcKernel<__nv_bfloat16, __nv_bfloat16, float, false, true, (cudnnKernelDataType_t)0>(cudnn::engines_precompiled::nchw2nhwc_params_t, __nv_bfloat16 const*, __nv_bfloat16*)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(10.464), 'mean_duration_us': np.float64(10.464), 'median_duration_us': np.float64(10.464), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(10.464), 'max_duration_us': np.float64(10.464)}, {'name': 'void cutlass__5x_cudnn::Kernel(cutlass_tensorop_bf16_s16816fprop_optimized_bf16_256x64_32x4_nhwc_align8::Params)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(98.176), 'mean_duration_us': np.float64(98.176), 'median_duration_us': np.float64(98.176), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(98.176), 'max_duration_us': np.float64(98.176)}]","[{'name': 'Memset (Device)', 'stream': 7, 'mean_duration_us': np.float64(2.11)}, {'name': 'void cudnn::engines_precompiled::nhwcToNchwKernel<__nv_bfloat16,...', 'stream': 7, 'mean_duration_us': np.float64(2.21)}, {'name': 'Memset (Device)', 'stream': 7, 'mean_duration_us': np.float64(2.27)}, {'name': 'void cudnn::engines_precompiled::nchwToNhwcKernel<__nv_bfloat16,...', 'stream': 7, 'mean_duration_us': np.float64(10.4)}, {'name': 'void cudnn::engines_precompiled::nchwToNhwcKernel<__nv_bfloat16,...', 'stream': 7, 'mean_duration_us': np.float64(10.46)}, {'name': 'void cutlass__5x_cudnn::Kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(122.848), 'mean_duration_us': np.float64(10.237333333333334), 'median_duration_us': np.float64(10.224), 'std_dev_duration_us': np.float64(0.21878045819699932), 'min_duration_us': np.float64(10.016), 'max_duration_us': np.float64(10.848)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(122.848), 'mean_duration_us': np.float64(10.237333333333334), 'median_duration_us': np.float64(10.224), 'std_dev_duration_us': np.float64(0.21878045819699932), 'min_duration_us': np.float64(10.016), 'max_duration_us': np.float64(10.848)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)', 'stream': 7, 'count': 25, 'total_duration_us': np.float64(111.20100000000001), 'mean_duration_us': np.float64(4.448040000000001), 'median_duration_us': np.float64(4.448), 'std_dev_duration_us': np.float64(0.10934037863479323), 'min_duration_us': np.float64(4.256), 'max_duration_us': np.float64(4.64)}]","[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_ke...', 'stream': 7, 'mean_duration_us': np.float64(4.45)}]",9.075338746907917,88.78079007903182 aten::gelu,elementwise,python3,CPU,thread 5617 (python3),"((4, 197, 3072), ())","('c10::BFloat16', '')","((605184, 3072, 1), ())","('', '')",12,5.04510498046875,5.04510498046875,5.055419921875,5.055419921875,0.039437883181483786,0.039437883181483786,4.990966796875,4.990966796875,5.1201171875,5.1201171875,60.541259765625,60.541259765625,107,"[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::GeluCUDAKernelImpl(at::TensorIteratorBase&, at::native::GeluType)::{lambda()#2}::operator()() const::{lambda()#4}::operator()() const::{lambda(c10::BFloat16)#1}, std::array >(int, at::native::GeluCUDAKernelImpl(at::TensorIteratorBase&, at::native::GeluType)::{lambda()#2}::operator()() const::{lambda()#4}::operator()() const::{lambda(c10::BFloat16)#1}, std::array)', 'stream': 7, 'count': 12, 'total_duration_us': np.float64(60.541999999999994), 'mean_duration_us': np.float64(5.045166666666666), 'median_duration_us': np.float64(5.0555), 'std_dev_duration_us': np.float64(0.037739972914080884), 'min_duration_us': np.float64(4.991), 'max_duration_us': np.float64(5.12)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Ge...', 'stream': 7, 'mean_duration_us': np.float64(5.05)}]",4.940897991876649,93.72168807090847 aten::add,elementwise,python3,CPU,thread 5617 (python3),"((4, 197, 768), (4, 197, 768), ())","('c10::BFloat16', 'c10::BFloat16', 'Scalar')","((151296, 768, 1), (151296, 768, 1), ())","('', '', '1')",24,2.2306620279947915,2.2306620279947915,2.2244873046875,2.2244873046875,0.03450867810431143,0.03450867810431143,2.176025390625,2.176025390625,2.303955078125,2.303955078125,53.535888671875,53.535888671875,90,"[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::CUDAFunctor_add, std::array >(int, at::native::CUDAFunctor_add, std::array)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(53.536), 'mean_duration_us': np.float64(2.2306666666666666), 'median_duration_us': np.float64(2.2245), 'std_dev_duration_us': np.float64(0.033798011118341686), 'min_duration_us': np.float64(2.176), 'max_duration_us': np.float64(2.304)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::CU...', 'stream': 7, 'mean_duration_us': np.float64(2.23)}]",4.369175102338875,98.09086317324734 diff --git a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv index 2d78a6869..5f57479ea 100644 --- a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv +++ b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv @@ -1,2 +1,3 @@ name,param: B,param: N_Q,param: H_Q,param: N_KV,param: H_KV,param: d_h_qk,param: d_h_v,param: dropout,param: causal,param: flash_impl,param: dtype_A_B,GFLOPS_first,Data Moved (MB)_first,FLOPS/Byte_first,TB/s_mean,TB/s_median,TB/s_std,TB/s_min,TB/s_max,TFLOPS/s_mean,TFLOPS/s_median,TFLOPS/s_std,TFLOPS/s_min,TFLOPS/s_max,process_name_first,process_label_first,thread_name_first,Compute Spec,kernel_details__summarize_kernel_stats,trunc_kernel_details,Input Dims_first,Input type_first,Input Strides_first,Concrete Inputs_first,Kernel Time (µs)_mean,Kernel Time (µs)_median,Kernel Time (µs)_std,Kernel Time (µs)_min,Kernel Time (µs)_max,Kernel Time (µs)_sum,name_count,UID_first aten::_scaled_dot_product_flash_attention,16,559,16,559,16,64,64,0.0,True,True,"('c10::BFloat16', 'c10::BFloat16')",10.239377408,69.875,139.75,0.6460767344679231,0.6439839021131257,0.012288781806023831,0.6123786199798398,0.6704724410436742,90.28922364189224,89.9967503203093,1.7173572573918277,85.57991214218261,93.69852363585345,python3,CPU,thread 25286 (python3),matrix_bf16,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel > >(at::TensorIteratorBase&, at::native::BinaryFunctor > const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast > >(at::TensorIteratorBase&, at::native::BinaryFunctor > const&)::{lambda(int)#1})",7,3913.711,145,26.991110344827586,23.392,32.672,8.614304613826809,11.216508107775581 elementwise,aten::copy_,"void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})",7,2766.267,96,28.815281249999998,26.913,30.048,6.088713903805581,7.927988610751288 -other,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,2722.705,24,113.44604166666666,109.28,119.647,5.992831418464297,7.803142007057015 +SDPA_fwd,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,2722.705,24,113.44604166666666,109.28,119.647,5.992831418464297,7.803142007057015 GEMM,aten::addmm,sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas,7,2266.292,72,31.476277777777778,30.912,32.256,4.988239967610992,6.4950842289037025 other,aten::cat,"void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)",7,2014.326,48,41.965125,41.376,42.464,4.433648206408521,5.772961751826632 elementwise,aten::add,"void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast >(at::TensorIteratorBase&, at::native::CUDAFunctor_add const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast >(at::TensorIteratorBase&, at::native::CUDAFunctor_add const&)::{lambda(int)#1})",7,1671.096,48,34.8145,34.271,35.615,3.6781790947127995,4.789281025827238 diff --git a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv index 1c0587fec..ea93a42fc 100644 --- a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv +++ b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv @@ -2,7 +2,7 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_ker aten::mul,6752.89990234375,6752.89990234375,6.75289990234375,6.75289990234375,220,['elementwise'],21.165327193718745,21.165327193718745 aten::mm,6668.681640625,6668.681640625,6.668681640625,6.668681640625,96,['GEMM'],20.901365474939983,42.06669266865873 aten::copy_,5101.934814453125,5101.934814453125,5.101934814453125,5.101934814453125,197,['elementwise'],15.99077747790201,58.057470146560746 -aten::_flash_attention_forward,2722.704345703125,2722.704345703125,2.722704345703125,2.722704345703125,24,['other'],8.53365652711153,66.59112667367228 +aten::_flash_attention_forward,2722.704345703125,2722.704345703125,2.722704345703125,2.722704345703125,24,['SDPA_fwd'],8.53365652711153,66.59112667367228 aten::add,2607.51513671875,2607.51513671875,2.60751513671875,2.60751513671875,145,['elementwise'],8.17262388445474,74.76375055812701 aten::addmm,2329.075927734375,2329.075927734375,2.329075927734375,2.329075927734375,72,['GEMM'],7.29992370424488,82.06367426237189 aten::cat,2017.782958984375,2017.782958984375,2.017782958984375,2.017782958984375,49,['other'],6.324251380949954,88.38792564332185 diff --git a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv index 8a3d282f4..1e19d5b37 100644 --- a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv @@ -1,5 +1,6 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) elementwise,735,17.360615234375,54.41263857506248,54.41263857506248 GEMM,169,9.001277587890625,28.212321826815455,82.62496040187793 -other,76,4.796135009765625,15.032322145291522,97.65728254716944 -reduce,50,0.747455322265625,2.3427174528305534,100.0 +SDPA_fwd,24,2.722704345703125,8.53365652711153,91.15861692898946 +other,52,2.0734306640625,6.498665618179992,97.65728254716946 +reduce,50,0.747455322265625,2.3427174528305534,100.00000000000001 diff --git a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv index 5cc0743f5..cc31a0fae 100644 --- a/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/h100/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv @@ -1,6 +1,6 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::mm,GEMM,python3,CPU,thread 25286 (python3),"((8944, 1024), (1024, 2816))","('c10::BFloat16', 'c10::BFloat16')","((1024, 1), (1, 1024))","('', '')",48,84.4844258626302,84.4844258626302,84.4635009765625,84.4635009765625,1.1997237297511993,1.1997237297511993,82.27197265625,82.27197265625,86.912109375,86.912109375,4055.25244140625,4055.25244140625,215,"[{'name': 'Memset (Device)', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(39.75), 'mean_duration_us': np.float64(0.828125), 'median_duration_us': np.float64(0.8), 'std_dev_duration_us': np.float64(0.07937112851450877), 'min_duration_us': np.float64(0.768), 'max_duration_us': np.float64(1.057)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(4015.5009999999997), 'mean_duration_us': np.float64(83.65627083333332), 'median_duration_us': np.float64(83.679), 'std_dev_duration_us': np.float64(1.1791844593401435), 'min_duration_us': np.float64(81.472), 'max_duration_us': np.float64(86.016)}]","[{'name': 'Memset (Device)', 'stream': 7, 'mean_duration_us': np.float64(0.83)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warp...', 'stream': 7, 'mean_duration_us': np.float64(83.66)}]",12.710205395714588,12.710205395714588 -aten::_flash_attention_forward,other,python3,CPU,thread 25286 (python3),"((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",24,113.44601440429688,113.44601440429688,113.7750244140625,113.7750244140625,2.173453067400901,2.173453067400901,109.280029296875,109.280029296875,119.64697265625,119.64697265625,2722.704345703125,2722.704345703125,164,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, true, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 24, 'total_duration_us': np.float64(2722.705), 'mean_duration_us': np.float64(113.44604166666666), 'median_duration_us': np.float64(113.775), 'std_dev_duration_us': np.float64(2.127694555913488), 'min_duration_us': np.float64(109.28), 'max_duration_us': np.float64(119.647)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 72, 'total_duration_us': np.float64(2108.187), 'mean_duration_us': np.float64(29.280375), 'median_duration_us': np.float64(29.312), 'std_dev_duration_us': np.float64(0.3210814381594732), 'min_duration_us': np.float64(28.672), 'max_duration_us': np.float64(30.048)}]","[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(29.28)}]",6.607603257237321,35.151388884308325 aten::cat,other,python3,CPU,thread 25286 (python3),"(((16, 16, 559, 32), (16, 16, 559, 32)), ())","('TensorList', 'Scalar')","(((286208, 32, 512, 1), (572416, 64, 1024, 1)), ())","('', '-1')",48,41.96514383951823,41.96514383951823,41.9678955078125,41.9678955078125,0.2933502450882166,0.2933502450882166,41.3759765625,41.3759765625,42.464111328125,42.464111328125,2014.326904296875,2014.326904296875,129,"[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)', 'stream': 7, 'count': 48, 'total_duration_us': np.float64(2014.326), 'mean_duration_us': np.float64(41.965125), 'median_duration_us': np.float64(41.968), 'std_dev_duration_us': np.float64(0.290296301575362), 'min_duration_us': np.float64(41.376), 'max_duration_us': np.float64(42.464)}]","[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(5328692.5649999995), 'mean_duration_us': np.float64(17762.308549999998), 'median_duration_us': np.float64(17749.3595), 'std_dev_duration_us': np.float64(234.40312254453335), 'min_duration_us': np.float64(17253.507), 'max_duration_us': np.float64(18602.391)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(5328692.5649999995), 'mean_duration_us': np.float64(17762.308549999998), 'median_duration_us': np.float64(17749.3595), 'std_dev_duration_us': np.float64(234.40312254453335), 'min_duration_us': np.float64(17253.507), 'max_duration_us': np.float64(18602.391)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(97402.64199999999), 'mean_duration_us': np.float64(324.6754733333333), 'median_duration_us': np.float64(324.733), 'std_dev_duration_us': np.float64(3.6598591506535088), 'min_duration_us': np.float64(316.99), 'max_duration_us': np.float64(337.213)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(97402.64199999999), 'mean_duration_us': np.float64(324.6754733333333), 'median_duration_us': np.float64(324.733), 'std_dev_duration_us': np.float64(3.6598591506535088), 'min_duration_us': np.float64(316.99), 'max_duration_us': np.float64(337.213)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel::Params)', 'stream': 7, 'count': 126, 'total_duration_us': np.float64(8900.819), 'mean_duration_us': np.float64(70.64142063492064), 'median_duration_us': np.float64(70.63900000000001), 'std_dev_duration_us': np.float64(0.41238584477528567), 'min_duration_us': np.float64(69.599), 'max_duration_us': np.float64(71.808)}]","[{'name': 'fmha_cutlassF_bf16_aligned_32x128_gmem_sm80(PyTorchMemEffAttenti...', 'stream': 7, 'mean_duration_us': np.float64(70.64)}]","[[1, 1, 1024, 384], [1, 1, 1024, 384], [1, 1, 1024, 384], [], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', 'Scalar', 'Scalar', 'Scalar', '']","[[1179648, 1179648, 1152, 1], [1179648, 1179648, 1152, 1], [1179648, 1179648, 1152, 1], [], [], [], [], []]","['', '', '', '', 'False', '0.', 'False', '']",70.64143492683532,70.6390380859375,0.4140213207002293,69.59912109375,71.80810546875,8900.82080078125,126,133631 diff --git a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/kernel_summary.csv b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/kernel_summary.csv index 1be4dbdda..aae8a2b79 100644 --- a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/kernel_summary.csv +++ b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/kernel_summary.csv @@ -1,5 +1,5 @@ Parent op category,Parent cpu_op,Kernel name,Kernel stream,Kernel duration (µs)_sum,Kernel duration (µs)_count,Kernel duration (µs)_mean,Kernel duration (µs)_min,Kernel duration (µs)_max,Percent of kernels time (%),Percent of total time (%) -other,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)",7,5426095.207,600,9043.492011666667,316.99,18602.391,47.31507844531886,45.82366108719575 +SDPA_fwd,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)",7,5426095.207,600,9043.492011666667,316.99,18602.391,47.31507844531886,45.82366108719575 GEMM,aten::addmm,sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas,7,710124.162,2110,336.5517355450237,49.983,1552.273,6.1922209524817715,5.997036116752548 elementwise,aten::copy_,"void at::native::unrolled_elementwise_kernel, 4, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1> >(int, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#7}::operator()() const::{lambda(float)#1}, std::array, TrivialOffsetCalculator<1, unsigned int>, TrivialOffsetCalculator<1, unsigned int>, at::native::memory::LoadWithCast<1>, at::native::memory::StoreWithCast<1>)",7,433032.332,4492,96.40078628673197,1.888,158.078,3.776004286011666,3.656980953037308 elementwise,aten::copy_,"void at::native::elementwise_kernel<128, 2, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#7}::operator()() const::{lambda(float)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#7}::operator()() const::{lambda(float)#1} const&)::{lambda(int)#1})",7,374296.156,1285,291.2810552529183,3.968,403.228,3.26382993797722,3.160950839317654 diff --git a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary.csv b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary.csv index 9f37decf1..89feba90a 100644 --- a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary.csv +++ b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary.csv @@ -1,5 +1,5 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) -aten::_flash_attention_forward,5426095.207763672,5426095.207763672,5426.095207763672,5426.095207763672,600,['other'],46.851463608237154,46.851463608237154 +aten::_flash_attention_forward,5426095.207763672,5426095.207763672,5426.095207763672,5426.095207763672,600,['SDPA_fwd'],46.851463608237154,46.851463608237154 aten::copy_,1625219.80859375,1625219.80859375,1625.21980859375,1625.21980859375,19010,['elementwise'],14.032913873086725,60.88437748132388 aten::addmm,1055648.8400878906,1055648.8400878906,1055.6488400878907,1055.6488400878907,3060,['GEMM'],9.114969664315867,69.99934714563975 aten::cudnn_convolution,928965.0529785156,928965.0529785156,928.9650529785156,928.9650529785156,4797,['CONV_fwd'],8.021122134140525,78.02046927978027 diff --git a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary_by_category.csv index a8082ab7e..e5690de6f 100644 --- a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_summary_by_category.csv @@ -1,7 +1,8 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) -other,1328,5470.613397216797,47.235854639578335,47.235854639578335 -elementwise,70504,3837.8522712402346,33.1378254776162,80.37368011719454 -GEMM,3492,1073.445296875,9.268632660557405,89.64231277775194 -CONV_fwd,4797,928.9650529785156,8.021122134140525,97.66343491189247 -NORM_fwd,910,142.23528198242187,1.2281264670903704,98.89156137898284 -reduce,5078,128.3736520996094,1.1084386210171528,100.0 +SDPA_fwd,600,5426.095207763672,46.85146360823717,46.85146360823717 +elementwise,70504,3837.8522712402346,33.137825477616204,79.98928908585337 +GEMM,3492,1073.445296875,9.268632660557406,89.25792174641077 +CONV_fwd,4797,928.9650529785156,8.021122134140526,97.2790438805513 +NORM_fwd,910,142.23528198242187,1.2281264670903707,98.50717034764168 +reduce,5078,128.3736520996094,1.108438621017153,99.61560896865883 +other,728,44.518189453125,0.3843910313411765,100.00000000000001 diff --git a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_unique_args.csv b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_unique_args.csv index 10caa3b91..a5cf74b01 100644 --- a/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/h100/Wan-AI_Wan2.1-T2V-1.3B-Diffusers__1016009_perf_report_csvs/ops_unique_args.csv @@ -1,5 +1,5 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) -aten::_flash_attention_forward,other,python3,CPU,thread 10586 (python3),"((1, 32256, 12, 128), (1, 32256, 12, 128), (1, 32256, 12, 128), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((1536, 1536, 128, 1), (1536, 1536, 128, 1), (49545216, 1536, 128, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '32256', '32256', '0.', 'False', 'False', '0.088388347648318433', '', '', '', '')",300,17762.308549804686,17762.308549804686,17749.359497070312,17749.359497070312,234.79477512060606,234.79477512060606,17253.507080078125,17253.507080078125,18602.39111328125,18602.39111328125,5328692.564941406,5328692.564941406,12022,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(5328692.5649999995), 'mean_duration_us': np.float64(17762.308549999998), 'median_duration_us': np.float64(17749.3595), 'std_dev_duration_us': np.float64(234.40312254453335), 'min_duration_us': np.float64(17253.507), 'max_duration_us': np.float64(18602.391)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(5328692.5649999995), 'mean_duration_us': np.float64(17762.308549999998), 'median_duration_us': np.float64(17749.3595), 'std_dev_duration_us': np.float64(234.40312254453335), 'min_duration_us': np.float64(17253.507), 'max_duration_us': np.float64(18602.391)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#7}::operator()() const::{lambda(float)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#7}::operator()() const::{lambda(float)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 910, 'total_duration_us': np.float64(365534.161), 'mean_duration_us': np.float64(401.68589120879125), 'median_duration_us': np.float64(401.66), 'std_dev_duration_us': np.float64(0.6290816210688958), 'min_duration_us': np.float64(400.22), 'max_duration_us': np.float64(403.228)}]","[{'name': 'void at::native::elementwise_kernel<128, 2, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(401.69)}]",3.156194249804388,52.33249721389973 aten::addmm,GEMM,python3,CPU,thread 10586 (python3),"((8960,), (32256, 1536), (1536, 8960), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (1536, 1), (1, 1536), (), ())","('', '', '', '1', '1')",300,1149.8357413736978,1149.8357413736978,1148.3740234375,1148.3740234375,23.43479543431236,23.43479543431236,1143.60595703125,1143.60595703125,1553.264892578125,1553.264892578125,344950.7224121094,344950.7224121094,12226,"[{'name': 'Memset (Device)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(284.666), 'mean_duration_us': np.float64(0.9488866666666667), 'median_duration_us': np.float64(0.96), 'std_dev_duration_us': np.float64(0.21814426531286327), 'min_duration_us': np.float64(0.703), 'max_duration_us': np.float64(1.536)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(344666.045), 'mean_duration_us': np.float64(1148.8868166666666), 'median_duration_us': np.float64(1147.4615), 'std_dev_duration_us': np.float64(23.393421280844223), 'min_duration_us': np.float64(1142.646), 'max_duration_us': np.float64(1552.273)}]","[{'name': 'Memset (Device)', 'stream': 7, 'mean_duration_us': np.float64(0.95)}, {'name': 'sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warp...', 'stream': 7, 'mean_duration_us': np.float64(1148.89)}]",2.9784671294750265,55.310964343374756 @@ -22,7 +22,7 @@ aten::copy_,elementwise,python3,CPU,thread 10586 (python3),"((1, 12, 32256, 128) aten::add,elementwise,python3,CPU,thread 10586 (python3),"((1, 32256, 1536), (1, 1, 1536), ())","('float', 'float', 'Scalar')","((49545216, 1536, 1), (9216, 1536, 1), ())","('', '', '1')",600,166.78735310872395,166.78735310872395,166.846923828125,166.846923828125,0.6603320925995817,0.6603320925995817,165.278076171875,165.278076171875,168.798095703125,168.798095703125,100072.41186523438,100072.41186523438,11905,"[{'name': 'void at::native::elementwise_kernel<128, 2, at::native::gpu_kernel_impl_nocast >(at::TensorIteratorBase&, at::native::CUDAFunctor_add const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast >(at::TensorIteratorBase&, at::native::CUDAFunctor_add const&)::{lambda(int)#1})', 'stream': 7, 'count': 600, 'total_duration_us': np.float64(100072.41399999999), 'mean_duration_us': np.float64(166.78735666666665), 'median_duration_us': np.float64(166.847), 'std_dev_duration_us': np.float64(0.6597810314457033), 'min_duration_us': np.float64(165.278), 'max_duration_us': np.float64(168.798)}]","[{'name': 'void at::native::elementwise_kernel<128, 2, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(166.79)}]",0.8640723730730295,81.7710341953647 aten::mul,elementwise,python3,CPU,thread 10586 (python3),"((1, 32256, 1536), (1, 1, 1536))","('float', 'float')","((49545216, 1536, 1), (1536, 1536, 1))","('', '')",600,166.3344844563802,166.3344844563802,166.429931640625,166.429931640625,0.6636278409602321,0.6636278409602321,164.383056640625,164.383056640625,168.669921875,168.669921875,99800.69067382812,99800.69067382812,11904,"[{'name': 'void at::native::elementwise_kernel<128, 2, at::native::gpu_kernel_impl_nocast > >(at::TensorIteratorBase&, at::native::BinaryFunctor > const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast > >(at::TensorIteratorBase&, at::native::BinaryFunctor > const&)::{lambda(int)#1})', 'stream': 7, 'count': 600, 'total_duration_us': np.float64(99800.693), 'mean_duration_us': np.float64(166.33448833333333), 'median_duration_us': np.float64(166.43), 'std_dev_duration_us': np.float64(0.6630732763909941), 'min_duration_us': np.float64(164.383), 'max_duration_us': np.float64(168.67)}]","[{'name': 'void at::native::elementwise_kernel<128, 2, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(166.33)}]",0.8617262042309233,82.63276039959563 aten::copy_,elementwise,python3,CPU,thread 10586 (python3),"((501, 3, 256, 256), (501, 3, 256, 256), ())","('c10::BFloat16', 'c10::BFloat16', 'Scalar')","((65536, 32833536, 256, 1), (65536, 32833536, 256, 1), ())","('', '', 'False')",1,98194.1259765625,98194.1259765625,98194.1259765625,98194.1259765625,,,98194.1259765625,98194.1259765625,98194.1259765625,98194.1259765625,98194.1259765625,98194.1259765625,417967,"[{'name': 'Memcpy DtoH (Device -> Pageable)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(98194.126), 'mean_duration_us': np.float64(98194.126), 'median_duration_us': np.float64(98194.126), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(98194.126), 'max_duration_us': np.float64(98194.126)}]","[{'name': 'Memcpy DtoH (Device -> Pageable)', 'stream': 7, 'mean_duration_us': np.float64(98194.13)}]",0.8478543673821112,83.48061476697774 -aten::_flash_attention_forward,other,python3,CPU,thread 10586 (python3),"((1, 32256, 12, 128), (1, 512, 12, 128), (1, 512, 12, 128), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((49545216, 1536, 128, 1), (786432, 1536, 128, 1), (786432, 1536, 128, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '32256', '512', '0.', 'False', 'False', '0.088388347648318433', '', '', '', '')",300,324.67547607421875,324.67547607421875,324.73291015625,324.73291015625,3.6659696541442255,3.6659696541442255,316.989990234375,316.989990234375,337.212890625,337.212890625,97402.64282226562,97402.64282226562,12168,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(97402.64199999999), 'mean_duration_us': np.float64(324.6754733333333), 'median_duration_us': np.float64(324.733), 'std_dev_duration_us': np.float64(3.6598591506535088), 'min_duration_us': np.float64(316.99), 'max_duration_us': np.float64(337.213)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, true, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 300, 'total_duration_us': np.float64(97402.64199999999), 'mean_duration_us': np.float64(324.6754733333333), 'median_duration_us': np.float64(324.733), 'std_dev_duration_us': np.float64(3.6598591506535088), 'min_duration_us': np.float64(316.99), 'max_duration_us': np.float64(337.213)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(int, float, float const*, float const*, float const*, float*, float*, float*)', 'stream': 7, 'count': 610, 'total_duration_us': np.float64(94543.31599999999), 'mean_duration_us': np.float64(154.98904262295082), 'median_duration_us': np.float64(154.943), 'std_dev_duration_us': np.float64(0.9304707718137512), 'min_duration_us': np.float64(152.286), 'max_duration_us': np.float64(158.335)}]","[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_ke...', 'stream': 7, 'mean_duration_us': np.float64(154.99)}]",0.8163315546386649,85.13796665037954 aten::copy_,elementwise,python3,CPU,thread 10586 (python3),"((1, 96, 6, 256, 256), (1, 96, 6, 256, 256), ())","('c10::BFloat16', 'c10::BFloat16', 'Scalar')","((38340864, 399384, 66564, 258, 1), (37748736, 393216, 65536, 256, 1), ())","('', '', 'False')",868,106.201416859429,106.201416859429,106.175048828125,106.175048828125,0.24370126975100181,0.24370126975100181,105.597900390625,105.597900390625,107.39111328125,107.39111328125,92182.82983398438,92182.82983398438,139371,"[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 868, 'total_duration_us': np.float64(92182.826), 'mean_duration_us': np.float64(106.20141244239632), 'median_duration_us': np.float64(106.175), 'std_dev_duration_us': np.float64(0.2435553248252166), 'min_duration_us': np.float64(105.598), 'max_duration_us': np.float64(107.391)}]","[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(106.2)}]",0.7959500030688262,85.93391665344836 aten::cat,elementwise,python3,CPU,thread 10586 (python3),"(((1, 96, 2, 256, 256), (1, 96, 4, 256, 256)), ())","('TensorList', 'Scalar')","(((12582912, 131072, 65536, 256, 1), (25165824, 262144, 65536, 256, 1)), ())","('', '2')",744,101.31975858954974,101.31975858954974,101.31103515625,101.31103515625,0.3151894966770076,0.3151894966770076,100.1591796875,100.1591796875,102.55908203125,102.55908203125,75381.900390625,75381.900390625,139405,"[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 744, 'total_duration_us': np.float64(26078.028000000002), 'mean_duration_us': np.float64(35.05111290322581), 'median_duration_us': np.float64(34.848), 'std_dev_duration_us': np.float64(0.41360644785166834), 'min_duration_us': np.float64(34.367), 'max_duration_us': np.float64(36.031)}, {'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1}>(int, at::native::gpu_kernel_impl_nocast(at::TensorIteratorBase&, at::native::direct_copy_kernel_cuda(at::TensorIteratorBase&)::{lambda()#3}::operator()() const::{lambda()#12}::operator()() const::{lambda(c10::BFloat16)#1} const&)::{lambda(int)#1})', 'stream': 7, 'count': 744, 'total_duration_us': np.float64(49303.873999999996), 'mean_duration_us': np.float64(66.26864784946235), 'median_duration_us': np.float64(66.368), 'std_dev_duration_us': np.float64(0.3622605655986425), 'min_duration_us': np.float64(65.343), 'max_duration_us': np.float64(67.167)}]","[{'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(35.05)}, {'name': 'void at::native::elementwise_kernel<128, 4, at::native::gpu_kern...', 'stream': 7, 'mean_duration_us': np.float64(66.27)}]",0.6508828591540166,86.58479951260239 diff --git a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/SDPA_fwd.csv b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/SDPA_fwd.csv index 0e9ba42cc..701155342 100644 --- a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/SDPA_fwd.csv +++ b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/SDPA_fwd.csv @@ -1,2 +1,3 @@ name,param: B,param: N_Q,param: H_Q,param: N_KV,param: H_KV,param: d_h_qk,param: d_h_v,param: dropout,param: causal,param: flash_impl,param: dtype_A_B,GFLOPS_first,Data Moved (MB)_first,FLOPS/Byte_first,TB/s_mean,TB/s_median,TB/s_std,TB/s_min,TB/s_max,TFLOPS/s_mean,TFLOPS/s_median,TFLOPS/s_std,TFLOPS/s_min,TFLOPS/s_max,process_name_first,process_label_first,thread_name_first,Compute Spec,kernel_details__summarize_kernel_stats,trunc_kernel_details,Input Dims_first,Input type_first,Input Strides_first,Concrete Inputs_first,Kernel Time (µs)_mean,Kernel Time (µs)_median,Kernel Time (µs)_std,Kernel Time (µs)_min,Kernel Time (µs)_max,Kernel Time (µs)_sum,name_count,UID_first aten::_scaled_dot_product_flash_attention,1,141,8,141,8,64,64,0.0,False,True,"('c10::BFloat16', 'c10::BFloat16')",0.040716288,0.55078125,70.5,0.061873964703565223,0.06332550208801799,0.003054615255047572,0.057294794032164316,0.06355006060606061,4.362114511601348,4.464447897205268,0.2153503754808538,4.039282979267584,4.480279272727273,python3,CPU,thread 19392 (python3),matrix_bf16,"[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 4, 'total_duration_us': np.float64(37.408), 'mean_duration_us': np.float64(9.352), 'median_duration_us': np.float64(9.12), 'std_dev_duration_us': np.float64(0.4205139712304459), 'min_duration_us': np.float64(9.088), 'max_duration_us': np.float64(10.08)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 4, 'total_duration_us': np.float64(37.408), 'mean_duration_us': np.float64(9.352), 'median_duration_us': np.float64(9.12), 'std_dev_duration_us': np.float64(0.4205139712304459), 'min_duration_us': np.float64(9.088), 'max_duration_us': np.float64(10.08)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_32x32_128x2_tn_align8::Params),7,98.718,17,5.8069411764705885,5.633,6.08,32.135079444135194,2.5261622829445454 GEMM,aten::addmm,void cutlass::Kernel2(cutlass_80_tensorop_bf16_s16816gemm_bf16_256x128_64x3_tn_align2::Params),7,41.28,1,41.28,41.28,41.28,13.437631226867449,1.0563420960711403 -other,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,37.408,4,9.352,9.088,10.08,12.177202251324069,0.9572588451993512 +SDPA_fwd,aten::_flash_attention_forward,"void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)",7,37.408,4,9.352,9.088,10.08,12.177202251324069,0.9572588451993512 GEMM,aten::addmm,sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize64x64x64_warpgroupsize1x1x1_execute_segment_k_off_kernel__5x_cublas,7,33.089,4,8.27225,8.16,8.48,10.77126404229208,0.8467370062233033 NORM_fwd,aten::native_layer_norm,"void at::native::(anonymous namespace)::vectorized_layer_norm_kernel(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)",7,30.912,10,3.0911999999999997,3.008,3.263,10.062598267561206,0.7910282672904818 GEMM,aten::addmm,void cutlass::Kernel2(cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x64_64x6_tn_align8::Params),7,24.511,4,6.12775,5.984,6.336,7.978919064964826,0.6272287092247995 diff --git a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv index 6693db80c..56bb2290c 100644 --- a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv +++ b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv @@ -1,6 +1,6 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::addmm,197.59814453125,197.59814453125,0.19759814453125,0.19759814453125,26,['GEMM'],63.857156495249946,63.857156495249946 -aten::_flash_attention_forward,37.408203125,37.408203125,0.037408203125,0.037408203125,4,['other'],12.089088623913868,75.94624511916382 +aten::_flash_attention_forward,37.408203125,37.408203125,0.037408203125,0.037408203125,4,['SDPA_fwd'],12.089088623913868,75.94624511916382 aten::native_layer_norm,30.912109375,30.912109375,0.030912109375,0.030912109375,10,['NORM_fwd'],9.989766911224603,85.93601203038843 aten::add,14.752685546875,14.752685546875,0.014752685546875,0.014752685546875,9,['elementwise'],4.7675779138858365,90.70358994427426 aten::index_select,9.823974609375,9.823974609375,0.009823974609375,0.009823974609375,3,['other'],3.1747822608577647,93.87837220513202 diff --git a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv index 4a48cf5d5..ed5e97c0b 100644 --- a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv @@ -1,6 +1,7 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) -GEMM,26,0.19759814453125,63.857156495249946,63.857156495249946 -other,8,0.04947216796875,15.987761320502395,79.84491781575234 -NORM_fwd,10,0.030912109375,9.989766911224603,89.83468472697695 -elementwise,16,0.027807373046875,8.98641926313871,98.82110399011566 -reduce,1,0.00364794921875,1.178896009884359,100.00000000000001 +GEMM,26,0.19759814453125,63.85715649524993,63.85715649524993 +SDPA_fwd,4,0.037408203125,12.089088623913867,75.94624511916379 +NORM_fwd,10,0.030912109375,9.989766911224601,85.9360120303884 +elementwise,16,0.027807373046875,8.986419263138709,94.92243129352711 +other,4,0.01206396484375,3.8986726965885232,98.82110399011563 +reduce,1,0.00364794921875,1.1788960098843588,99.99999999999999 diff --git a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv index 388f33856..e89927304 100644 --- a/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/h100/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv @@ -1,7 +1,7 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::addmm,GEMM,python3,CPU,thread 19392 (python3),"((512,), (141, 512), (512, 512), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (512, 1), (1, 512), (), ())","('', '', '', '1', '1')",17,5.806942210477941,5.806942210477941,5.7919921875,5.7919921875,0.11484330395890131,0.11484330395890131,5.633056640625,5.633056640625,6.080078125,6.080078125,98.718017578125,98.718017578125,48,"[{'name': 'void cutlass::Kernel2(cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_32x32_128x2_tn_align8::Params)', 'stream': 7, 'count': 17, 'total_duration_us': np.float64(98.718), 'mean_duration_us': np.float64(5.8069411764705885), 'median_duration_us': np.float64(5.792), 'std_dev_duration_us': np.float64(0.11144002374700734), 'min_duration_us': np.float64(5.633), 'max_duration_us': np.float64(6.08)}]","[{'name': 'void cutlass::Kernel2(cutlass_80_tensorop_bf16_s16816gemm_bf16_256x128_64x3_tn_align2::Params)', 'stream': 7, 'count': 1, 'total_duration_us': np.float64(41.28), 'mean_duration_us': np.float64(41.28), 'median_duration_us': np.float64(41.28), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(41.28), 'max_duration_us': np.float64(41.28)}]","[{'name': 'void cutlass::Kernel2 >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 4, 'total_duration_us': np.float64(37.408), 'mean_duration_us': np.float64(9.352), 'median_duration_us': np.float64(9.12), 'std_dev_duration_us': np.float64(0.4205139712304459), 'min_duration_us': np.float64(9.088), 'max_duration_us': np.float64(10.08)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel >, false, false, false, false, false, true, false>(pytorch_flash::Flash_fwd_params)', 'stream': 7, 'count': 4, 'total_duration_us': np.float64(37.408), 'mean_duration_us': np.float64(9.352), 'median_duration_us': np.float64(9.12), 'std_dev_duration_us': np.float64(0.4205139712304459), 'min_duration_us': np.float64(9.088), 'max_duration_us': np.float64(10.08)}]","[{'name': 'void pytorch_flash::flash_fwd_kernel(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)', 'stream': 7, 'count': 10, 'total_duration_us': np.float64(30.911999999999995), 'mean_duration_us': np.float64(3.0911999999999997), 'median_duration_us': np.float64(3.056), 'std_dev_duration_us': np.float64(0.08591833331716812), 'min_duration_us': np.float64(3.008), 'max_duration_us': np.float64(3.263)}]","[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_ke...', 'stream': 7, 'mean_duration_us': np.float64(3.09)}]",9.989766911224601,78.01487545534088 aten::addmm,GEMM,python3,CPU,thread 19392 (python3),"((2048,), (141, 512), (512, 2048), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (512, 1), (1, 512), (), ())","('', '', '', '1', '1')",4,6.12774658203125,6.12774658203125,6.0955810546875,6.0955810546875,0.1536288548372083,0.1536288548372083,5.98388671875,5.98388671875,6.3359375,6.3359375,24.510986328125,24.510986328125,123,"[{'name': 'void cutlass::Kernel2(cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x64_64x6_tn_align8::Params)', 'stream': 7, 'count': 4, 'total_duration_us': np.float64(24.510999999999996), 'mean_duration_us': np.float64(6.127749999999999), 'median_duration_us': np.float64(6.0954999999999995), 'std_dev_duration_us': np.float64(0.13305708361451504), 'min_duration_us': np.float64(5.984), 'max_duration_us': np.float64(6.336)}]","[{'name': 'void cutlass::Kernel2(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)",0,139.695,25,5.5878,4.81,6.293,8.139306566855115,2.0337468116894066 diff --git a/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv b/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv index 8a246b3f9..6b85d5cf4 100644 --- a/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv +++ b/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary.csv @@ -1,6 +1,6 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::addmm,942.326171875,942.326171875,0.942326171875,0.942326171875,73,['GEMM'],54.90455137724823,54.90455137724823 -aten::_flash_attention_forward,374.68701171875,374.68701171875,0.37468701171875,0.37468701171875,12,['other'],21.831105724641393,76.73565710188963 +aten::_flash_attention_forward,374.68701171875,374.68701171875,0.37468701171875,0.37468701171875,12,['SDPA_fwd'],21.831105724641393,76.73565710188963 aten::native_layer_norm,139.6953125,139.6953125,0.1396953125,0.1396953125,25,['NORM_fwd'],8.13933507445277,84.8749921763424 aten::add,130.03369140625,130.03369140625,0.13003369140625,0.13003369140625,25,['elementwise'],7.576401572697428,92.45139374903982 aten::gelu,65.6572265625,65.6572265625,0.0656572265625,0.0656572265625,12,['elementwise'],3.825512520697131,96.27690626973695 diff --git a/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv index 727b55f20..4674d8c85 100644 --- a/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_summary_by_category.csv @@ -1,6 +1,7 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) GEMM,73,0.942326171875,54.90455137724823,54.90455137724823 -other,13,0.3814609375,22.22578791344474,77.13033929069297 -elementwise,39,0.206474609375,12.030224923043658,89.16056421373663 -NORM_fwd,25,0.1396953125,8.13933507445277,97.29989928818941 -CONV_fwd,1,0.046341796875,2.7001007118105935,100.0 +SDPA_fwd,12,0.37468701171875,21.831105724641393,76.73565710188963 +elementwise,39,0.206474609375,12.030224923043658,88.76588202493329 +NORM_fwd,25,0.1396953125,8.13933507445277,96.90521709938606 +CONV_fwd,1,0.046341796875,2.7001007118105935,99.60531781119666 +other,1,0.00677392578125,0.3946821888033502,100.0 diff --git a/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv b/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv index 27b5a62ad..e2ba20ed2 100644 --- a/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/mi300/Falconsai_nsfw_image_detection__1016002_perf_report_csvs/ops_unique_args.csv @@ -1,6 +1,6 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::addmm,GEMM,python3,CPU,thread 950 (python3),"((768,), (788, 768), (768, 768), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (768, 1), (1, 768), (), ())","('', '', '', '1', '1')",48,8.525146484375,8.525146484375,8.5380859375,8.5380859375,0.532247235181306,0.532247235181306,7.2958984375,7.2958984375,9.7412109375,9.7412109375,409.20703125,409.20703125,33,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x128_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(409.207), 'mean_duration_us': np.float64(8.525145833333333), 'median_duration_us': np.float64(8.538), 'std_dev_duration_us': np.float64(0.526665611401237), 'min_duration_us': np.float64(7.296), 'max_duration_us': np.float64(9.741)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x128_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(8.53)}]",23.842411621118757,23.842411621118757 -aten::_flash_attention_forward,other,python3,CPU,thread 950 (python3),"((4, 197, 12, 64), (4, 197, 12, 64), (4, 197, 12, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((151296, 768, 64, 1), (151296, 768, 64, 1), (151296, 768, 64, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '197', '197', '0.', 'False', 'False', '0.125', '', '', '', '')",12,31.223917643229168,31.223917643229168,31.2099609375,31.2099609375,0.6886496730633799,0.6886496730633799,30.42919921875,30.42919921875,32.875,32.875,374.68701171875,374.68701171875,71,"[{'name': 'attn_fwd', 'stream': 0, 'count': 12, 'total_duration_us': np.float64(374.687), 'mean_duration_us': np.float64(31.223916666666668), 'median_duration_us': np.float64(31.21), 'std_dev_duration_us': np.float64(0.6592922288754476), 'min_duration_us': np.float64(30.429), 'max_duration_us': np.float64(32.875)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(31.22)}]",21.831105724641393,45.673517345760146 +aten::_flash_attention_forward,SDPA_fwd,python3,CPU,thread 950 (python3),"((4, 197, 12, 64), (4, 197, 12, 64), (4, 197, 12, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((151296, 768, 64, 1), (151296, 768, 64, 1), (151296, 768, 64, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '197', '197', '0.', 'False', 'False', '0.125', '', '', '', '')",12,31.223917643229168,31.223917643229168,31.2099609375,31.2099609375,0.6886496730633799,0.6886496730633799,30.42919921875,30.42919921875,32.875,32.875,374.68701171875,374.68701171875,71,"[{'name': 'attn_fwd', 'stream': 0, 'count': 12, 'total_duration_us': np.float64(374.687), 'mean_duration_us': np.float64(31.223916666666668), 'median_duration_us': np.float64(31.21), 'std_dev_duration_us': np.float64(0.6592922288754476), 'min_duration_us': np.float64(30.429), 'max_duration_us': np.float64(32.875)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(31.22)}]",21.831105724641393,45.673517345760146 aten::addmm,GEMM,python3,CPU,thread 950 (python3),"((768,), (788, 3072), (3072, 768), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (3072, 1), (1, 3072), (), ())","('', '', '', '1', '1')",12,22.871622721354168,22.871622721354168,22.8115234375,22.8115234375,1.324204222903877,1.324204222903877,21.248046875,21.248046875,25.13720703125,25.13720703125,274.45947265625,274.45947265625,127,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB1024_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_2_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB2_WSGRA0_WSGRB0_WS64_WG64_4_1', 'stream': 0, 'count': 12, 'total_duration_us': np.float64(274.459), 'mean_duration_us': np.float64(22.871583333333334), 'median_duration_us': np.float64(22.811500000000002), 'std_dev_duration_us': np.float64(1.2678273711572705), 'min_duration_us': np.float64(21.248), 'max_duration_us': np.float64(25.137)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT64x32x256_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(22.87)}]",15.991356992073925,61.66487433783407 aten::addmm,GEMM,python3,CPU,thread 950 (python3),"((3072,), (788, 768), (768, 3072), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (768, 1), (1, 768), (), ())","('', '', '', '1', '1')",12,20.930216471354168,20.930216471354168,20.967529296875,20.967529296875,0.9688740862432101,0.9688740862432101,19.40380859375,19.40380859375,22.4501953125,22.4501953125,251.16259765625,251.16259765625,117,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT64x160x128_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_5_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 12, 'total_duration_us': np.float64(251.16200000000003), 'mean_duration_us': np.float64(20.93016666666667), 'median_duration_us': np.float64(20.9675), 'std_dev_duration_us': np.float64(0.9275455813896989), 'min_duration_us': np.float64(19.404), 'max_duration_us': np.float64(22.45)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT64x160x128_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(20.93)}]",14.633966622854183,76.29884096068825 aten::native_layer_norm,NORM_fwd,python3,CPU,thread 950 (python3),"((4, 197, 768), (), (768,), (768,), ())","('c10::BFloat16', 'ScalarList', 'c10::BFloat16', 'c10::BFloat16', 'Scalar')","((151296, 768, 1), (), (1,), (1,), ())","('', '[768]', '', '', '9.9999999999999998e-13')",25,5.5878125,5.5878125,5.73193359375,5.73193359375,0.44789382994639776,0.44789382994639776,4.81005859375,4.81005859375,6.29296875,6.29296875,139.6953125,139.6953125,21,"[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_kernel(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)', 'stream': 0, 'count': 25, 'total_duration_us': np.float64(139.695), 'mean_duration_us': np.float64(5.5878), 'median_duration_us': np.float64(5.732), 'std_dev_duration_us': np.float64(0.4388615271358382), 'min_duration_us': np.float64(4.81), 'max_duration_us': np.float64(6.293)}]","[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_ke...', 'stream': 0, 'mean_duration_us': np.float64(5.59)}]",8.13933507445277,84.43817603514103 diff --git a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv index 6bd6955b3..2341b9bbb 100644 --- a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv +++ b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/SDPA_fwd.csv @@ -1,2 +1,3 @@ name,param: B,param: N_Q,param: H_Q,param: N_KV,param: H_KV,param: d_h_qk,param: d_h_v,param: dropout,param: causal,param: flash_impl,param: dtype_A_B,GFLOPS_first,Data Moved (MB)_first,FLOPS/Byte_first,TB/s_mean,TB/s_median,TB/s_std,TB/s_min,TB/s_max,TFLOPS/s_mean,TFLOPS/s_median,TFLOPS/s_std,TFLOPS/s_min,TFLOPS/s_max,process_name_first,process_label_first,thread_name_first,Compute Spec,kernel_details__summarize_kernel_stats,trunc_kernel_details,Input Dims_first,Input type_first,Input Strides_first,Concrete Inputs_first,Kernel Time (µs)_mean,Kernel Time (µs)_median,Kernel Time (µs)_std,Kernel Time (µs)_min,Kernel Time (µs)_max,Kernel Time (µs)_sum,name_count,UID_first aten::_scaled_dot_product_flash_attention,16,559,16,559,16,64,64,0.0,True,True,"('c10::BFloat16', 'c10::BFloat16')",10.239377408,69.875,139.75,0.43959076839825145,0.43379994577448006,0.03016579398980997,0.4117933009618683,0.525005230283713,61.432809883655636,60.62354242198358,4.215669710075947,57.54811380942109,73.36948093214889,python3,CPU,thread 23232 (python3),matrix_bf16,"[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor, std::array >(int, at::native::FillFunctor, std::array)', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(107.545), 'mean_duration_us': np.float64(4.481041666666667), 'median_duration_us': np.float64(4.69), 'std_dev_duration_us': np.float64(0.9520142890719772), 'min_duration_us': np.float64(0.681), 'max_duration_us': np.float64(5.411)}, {'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Fi...', 'stream': 0, 'mean_duration_us': np.float64(4.48)}, {'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]","[[16, 16, 559, 64], [16, 16, 559, 64], [16, 16, 559, 64], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar', 'Scalar', 'Scalar']","[[572416, 35776, 64, 1], [572416, 35776, 64, 1], [572416, 35776, 64, 1], [], [], [], []]","['', '', '', '0.', 'True', 'False', '0.125']",167.3476359049479,168.906494140625,10.235203983555547,139.55908203125,177.92724609375,4016.34326171875,24,158 +aten::_flash_attention_forward,16,559,16,559,16,64,64,0.0,True,True,,10.239377408,69.875,139.75,0.43959076839825145,0.43379994577448006,0.03016579398980997,0.4117933009618683,0.525005230283713,61.432809883655636,60.62354242198358,4.215669710075947,57.54811380942109,73.36948093214889,python3,CPU,thread 23232 (python3),,"[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor, std::array >(int, at::native::FillFunctor, std::array)', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(107.545), 'mean_duration_us': np.float64(4.481041666666667), 'median_duration_us': np.float64(4.69), 'std_dev_duration_us': np.float64(0.9520142890719772), 'min_duration_us': np.float64(0.681), 'max_duration_us': np.float64(5.411)}, {'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'void at::native::vectorized_elementwise_kernel<4, at::native::Fi...', 'stream': 0, 'mean_duration_us': np.float64(4.48)}, {'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]","[[16, 559, 16, 64], [16, 559, 16, 64], [16, 559, 16, 64], [], [], [], [], [], [], [], [], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '']","[[572416, 64, 35776, 1], [572416, 64, 35776, 1], [572416, 64, 35776, 1], [], [], [], [], [], [], [], [], [], [], [], []]","['', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '']",167.3476359049479,168.906494140625,10.235203983555547,139.55908203125,177.92724609375,4016.34326171875,24,165 diff --git a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/kernel_summary.csv b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/kernel_summary.csv index b0fd44a33..6d952de21 100644 --- a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/kernel_summary.csv +++ b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/kernel_summary.csv @@ -1,6 +1,6 @@ Parent op category,Parent cpu_op,Kernel name,Kernel stream,Kernel duration (µs)_sum,Kernel duration (µs)_count,Kernel duration (µs)_mean,Kernel duration (µs)_min,Kernel duration (µs)_max,Percent of kernels time (%),Percent of total time (%) GEMM,aten::mm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_7_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1,0,5308.273,48,110.58902083333334,105.0,118.551,12.327035025417512,16.021911706710675 -other,aten::_flash_attention_forward,attn_fwd,0,3908.798,24,162.86658333333332,138.557,173.077,9.07713108449432,11.797889150646034 +SDPA_fwd,aten::_flash_attention_forward,attn_fwd,0,3908.798,24,162.86658333333332,138.557,173.077,9.07713108449432,11.797889150646034 GEMM,aten::mm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0,3355.344,48,69.903,39.209,105.883,7.791883162438047,10.127404018904345 GEMM,aten::addmm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0,2954.763,72,41.038375,39.57,45.023,6.861641628606464,8.918334060862271 ,,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1,0,2787.634,52,53.608346153846156,39.811,102.314,6.473529518177516,8.41389013312328 diff --git a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv index 213411270..0c046937b 100644 --- a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv +++ b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary.csv @@ -1,7 +1,7 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::mm,8663.615234375,8663.615234375,8.663615234375,8.663615234375,96,['GEMM'],29.12734092062252,29.12734092062252 aten::mul,4541.28857421875,4541.28857421875,4.54128857421875,4.54128857421875,220,['elementwise'],15.267951881722706,44.395292802345224 -aten::_flash_attention_forward,3908.79833984375,4016.34326171875,3.90879833984375,4.01634326171875,24,['other'],13.141500257635307,57.53679305998053 +aten::_flash_attention_forward,3908.79833984375,4016.34326171875,3.90879833984375,4.01634326171875,24,['SDPA_fwd'],13.141500257635307,57.53679305998053 aten::copy_,3190.6943359375,3190.6943359375,3.1906943359375,3.1906943359375,197,['elementwise'],10.727212506806964,68.26400556678749 aten::addmm,2954.7626953125,2954.7626953125,2.9547626953125,2.9547626953125,72,['GEMM'],9.934003073500232,78.19800864028772 aten::add,1895.603515625,1895.603515625,1.895603515625,1.895603515625,145,['elementwise'],6.373077330450393,84.57108597073812 diff --git a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv index b0c6b47d4..2cd86061c 100644 --- a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_summary_by_category.csv @@ -1,5 +1,6 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) elementwise,759,12.18791064453125,40.97613050090121,40.97613050090121 GEMM,169,11.6457197265625,39.153267956281255,80.12939845718248 -other,76,5.288919921875,17.78151147053321,97.9109099277157 -reduce,50,0.62137744140625,2.0890900722843178,100.00000000000001 +SDPA_fwd,24,3.90879833984375,13.141500257635307,93.27089871481778 +other,52,1.38012158203125,4.640011212897906,97.91090992771568 +reduce,50,0.62137744140625,2.0890900722843178,100.0 diff --git a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv index 5fc72f823..df0d0b388 100644 --- a/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/mi300/Qwen_Qwen1.5-0.5B-Chat__1016005_perf_report_csvs/ops_unique_args.csv @@ -1,6 +1,6 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::mm,GEMM,python3,CPU,thread 23232 (python3),"((8944, 1024), (1024, 2816))","('c10::BFloat16', 'c10::BFloat16')","((1024, 1), (1, 1024))","('', '')",48,110.58902994791667,110.58902994791667,110.452880859375,110.452880859375,3.280588398059727,3.280588398059727,105.0,105.0,118.55078125,118.55078125,5308.2734375,5308.2734375,231,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_7_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB1_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(5308.273), 'mean_duration_us': np.float64(110.58902083333334), 'median_duration_us': np.float64(110.453), 'std_dev_duration_us': np.float64(3.2462516941953186), 'min_duration_us': np.float64(105.0), 'max_duration_us': np.float64(118.551)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT128x224x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(110.59)}]",17.846578585400604,17.846578585400604 -aten::_flash_attention_forward,other,python3,CPU,thread 23232 (python3),"((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",24,162.8665974934896,167.3476359049479,164.237060546875,168.906494140625,9.452351716739908,10.235203983555547,138.55712890625,139.55908203125,173.0771484375,177.92724609375,3908.79833984375,4016.34326171875,165,"[{'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]",13.141500257635306,30.988078843035908 +aten::_flash_attention_forward,SDPA_fwd,python3,CPU,thread 23232 (python3),"((16, 559, 16, 64), (16, 559, 16, 64), (16, 559, 16, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((572416, 64, 35776, 1), (572416, 64, 35776, 1), (572416, 64, 35776, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '559', '559', '0.', 'True', 'False', '0.125', '', '', '', '')",24,162.8665974934896,167.3476359049479,164.237060546875,168.906494140625,9.452351716739908,10.235203983555547,138.55712890625,139.55908203125,173.0771484375,177.92724609375,3908.79833984375,4016.34326171875,165,"[{'name': 'attn_fwd', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(3908.798), 'mean_duration_us': np.float64(162.86658333333332), 'median_duration_us': np.float64(164.23700000000002), 'std_dev_duration_us': np.float64(9.253315184825505), 'min_duration_us': np.float64(138.557), 'max_duration_us': np.float64(173.077)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(162.87)}]",13.141500257635306,30.988078843035908 aten::addmm,GEMM,python3,CPU,thread 23232 (python3),"((1024,), (8944, 1024), (1024, 1024), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (1024, 1), (1, 1024), (), ())","('', '', '', '1', '1')",72,41.038370768229164,41.038370768229164,40.71240234375,40.71240234375,1.1213617423390887,1.1213617423390887,39.56982421875,39.56982421875,45.02294921875,45.02294921875,2954.7626953125,2954.7626953125,90,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 72, 'total_duration_us': np.float64(2954.763), 'mean_duration_us': np.float64(41.038375), 'median_duration_us': np.float64(40.7125), 'std_dev_duration_us': np.float64(1.1135532272552386), 'min_duration_us': np.float64(39.57), 'max_duration_us': np.float64(45.023)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(41.04)}]",9.93400307350023,40.922081916536136 aten::mm,GEMM,python3,CPU,thread 23232 (python3),"((8944, 2816), (2816, 1024))","('c10::BFloat16', 'c10::BFloat16')","((2816, 1), (1, 2816))","('', '')",24,98.94974772135417,98.94974772135417,98.324951171875,98.324951171875,3.3849417047301227,3.3849417047301227,93.65380859375,93.65380859375,105.8828125,105.8828125,2374.7939453125,2374.7939453125,253,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA1024_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_4_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB4_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 24, 'total_duration_us': np.float64(2374.795), 'mean_duration_us': np.float64(98.94979166666667), 'median_duration_us': np.float64(98.32499999999999), 'std_dev_duration_us': np.float64(3.3136632546066833), 'min_duration_us': np.float64(93.654), 'max_duration_us': np.float64(105.883)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT256x128x64_MI16x...', 'stream': 0, 'mean_duration_us': np.float64(98.95)}]",7.9841302955021485,48.906212212038284 aten::cat,other,python3,CPU,thread 23232 (python3),"(((16, 16, 559, 32), (16, 16, 559, 32)), ())","('TensorList', 'Scalar')","(((286208, 32, 512, 1), (572416, 64, 1024, 1)), ())","('', '-1')",48,28.119517008463543,28.119517008463543,27.90283203125,27.90283203125,0.8371311008823814,0.8371311008823814,26.81982421875,26.81982421875,30.34814453125,30.34814453125,1349.73681640625,1349.73681640625,130,"[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy, unsigned int, 4, 64, 64>(at::native::(anonymous namespace)::OpaqueType<2u>*, at::native::(anonymous namespace)::CatArrInputTensorMetadata, unsigned int, 64, 64>, at::native::(anonymous namespace)::TensorSizeStride, int, unsigned int)', 'stream': 0, 'count': 48, 'total_duration_us': np.float64(1349.738), 'mean_duration_us': np.float64(28.119541666666667), 'median_duration_us': np.float64(27.903), 'std_dev_duration_us': np.float64(0.8283528575415325), 'min_duration_us': np.float64(26.82), 'max_duration_us': np.float64(30.348)}]","[{'name': 'void at::native::(anonymous namespace)::CatArrayBatchedCopy(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)",0,39.519,10,3.9518999999999997,3.728,4.209,10.176287003275446,1.1304139147700363 GEMM,aten::addmm,Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT32x64x128_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1,0,33.794000000000004,4,8.448500000000001,7.095,12.147,8.702078569515688,0.9666542128024144 diff --git a/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv b/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv index 357ff9caa..3df68e2a7 100644 --- a/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv +++ b/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary.csv @@ -1,6 +1,6 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::addmm,174.130859375,174.130859375,0.174130859375,0.174130859375,26,['GEMM'],44.56340682259236,44.56340682259236 -aten::_flash_attention_forward,113.65576171875,113.65576171875,0.11365576171875,0.11365576171875,4,['other'],29.08667633860792,73.65008316120029 +aten::_flash_attention_forward,113.65576171875,113.65576171875,0.11365576171875,0.11365576171875,4,['SDPA_fwd'],29.08667633860792,73.65008316120029 aten::native_layer_norm,39.51904296875,39.51904296875,0.03951904296875,0.03951904296875,10,['NORM_fwd'],10.11367654979113,83.76375971099142 aten::gelu,21.7236328125,21.7236328125,0.0217236328125,0.0217236328125,5,['elementwise'],5.559491810714863,89.32325152170628 aten::add,19.3154296875,19.3154296875,0.0193154296875,0.0193154296875,9,['elementwise'],4.943186717200685,94.26643823890697 diff --git a/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv index 8df6f25c0..6a6e51da9 100644 --- a/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_summary_by_category.csv @@ -1,6 +1,7 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) -GEMM,26,0.174130859375,44.56340682259236,44.56340682259236 -other,8,0.12511767578125,32.0199986754189,76.58340549801126 -elementwise,16,0.04765185546875,12.195018325454575,88.77842382346583 -NORM_fwd,10,0.03951904296875,10.11367654979113,98.89210037325697 -reduce,1,0.0043291015625,1.1078996267430425,100.00000000000001 +GEMM,26,0.174130859375,44.563406822592356,44.563406822592356 +SDPA_fwd,4,0.11365576171875,29.086676338607916,73.65008316120027 +elementwise,16,0.04765185546875,12.195018325454573,85.84510148665484 +NORM_fwd,10,0.03951904296875,10.113676549791128,95.95877803644596 +other,4,0.0114619140625,2.933322336810983,98.89210037325695 +reduce,1,0.0043291015625,1.1078996267430423,100.0 diff --git a/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv b/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv index 94c10593d..383deab17 100644 --- a/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/mi300/gaunernst_bert-small-uncased__1016001_perf_report_csvs/ops_unique_args.csv @@ -1,5 +1,5 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) -aten::_flash_attention_forward,other,python3,CPU,thread 6240 (python3),"((1, 141, 8, 64), (1, 141, 8, 64), (1, 141, 8, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((72192, 512, 64, 1), (72192, 512, 64, 1), (72192, 512, 64, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '141', '141', '0.', 'False', 'False', '0.125', '', '', '', '')",4,28.4139404296875,28.4139404296875,27.3212890625,27.3212890625,2.4175726751563076,2.4175726751563076,26.98095703125,26.98095703125,32.0322265625,32.0322265625,113.65576171875,113.65576171875,86,"[{'name': 'attn_fwd', 'stream': 0, 'count': 4, 'total_duration_us': np.float64(113.656), 'mean_duration_us': np.float64(28.414), 'median_duration_us': np.float64(27.3215), 'std_dev_duration_us': np.float64(2.0935189275475854), 'min_duration_us': np.float64(26.981), 'max_duration_us': np.float64(32.032)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(28.41)}]",29.086676338607916,29.086676338607916 +aten::_flash_attention_forward,SDPA_fwd,python3,CPU,thread 6240 (python3),"((1, 141, 8, 64), (1, 141, 8, 64), (1, 141, 8, 64), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((72192, 512, 64, 1), (72192, 512, 64, 1), (72192, 512, 64, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '141', '141', '0.', 'False', 'False', '0.125', '', '', '', '')",4,28.4139404296875,28.4139404296875,27.3212890625,27.3212890625,2.4175726751563076,2.4175726751563076,26.98095703125,26.98095703125,32.0322265625,32.0322265625,113.65576171875,113.65576171875,86,"[{'name': 'attn_fwd', 'stream': 0, 'count': 4, 'total_duration_us': np.float64(113.656), 'mean_duration_us': np.float64(28.414), 'median_duration_us': np.float64(27.3215), 'std_dev_duration_us': np.float64(2.0935189275475854), 'min_duration_us': np.float64(26.981), 'max_duration_us': np.float64(32.032)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(28.41)}]",29.086676338607916,29.086676338607916 aten::addmm,GEMM,python3,CPU,thread 6240 (python3),"((512,), (141, 512), (512, 512), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (512, 1), (1, 512), (), ())","('', '', '', '1', '1')",17,5.043313419117647,5.043313419117647,5.009765625,5.009765625,0.22600854465708484,0.22600854465708484,4.72900390625,4.72900390625,5.4912109375,5.4912109375,85.736328125,85.736328125,48,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x128_MI16x16x1_SN_LDSB0_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA256_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT1_1_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS64_WG16_4_4', 'stream': 0, 'count': 17, 'total_duration_us': np.float64(85.736), 'mean_duration_us': np.float64(5.043294117647059), 'median_duration_us': np.float64(5.01), 'std_dev_duration_us': np.float64(0.2192558979569931), 'min_duration_us': np.float64(4.729), 'max_duration_us': np.float64(5.491)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT16x16x128_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(5.04)}]",21.94156098133965,51.028237319947564 aten::native_layer_norm,NORM_fwd,python3,CPU,thread 6240 (python3),"((1, 141, 512), (), (512,), (512,), ())","('c10::BFloat16', 'ScalarList', 'c10::BFloat16', 'c10::BFloat16', 'Scalar')","((72192, 512, 1), (), (1,), (1,), ())","('', '[512]', '', '', '9.9999999999999998e-13')",10,3.951904296875,3.951904296875,3.947998046875,3.947998046875,0.18037338930055438,0.18037338930055438,3.72802734375,3.72802734375,4.208984375,4.208984375,39.51904296875,39.51904296875,29,"[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_kernel(int, float, c10::BFloat16 const*, c10::BFloat16 const*, c10::BFloat16 const*, float*, float*, c10::BFloat16*)', 'stream': 0, 'count': 10, 'total_duration_us': np.float64(39.519000000000005), 'mean_duration_us': np.float64(3.9519000000000006), 'median_duration_us': np.float64(3.948), 'std_dev_duration_us': np.float64(0.17117561158062206), 'min_duration_us': np.float64(3.728), 'max_duration_us': np.float64(4.209)}]","[{'name': 'void at::native::(anonymous namespace)::vectorized_layer_norm_ke...', 'stream': 0, 'mean_duration_us': np.float64(3.95)}]",10.113676549791128,61.141913869738694 aten::addmm,GEMM,python3,CPU,thread 6240 (python3),"((2048,), (141, 512), (512, 2048), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar')","((1,), (512, 1), (1, 512), (), ())","('', '', '', '1', '1')",4,8.4486083984375,8.4486083984375,7.276123046875,7.276123046875,2.469693365812908,2.469693365812908,7.09521484375,7.09521484375,12.14697265625,12.14697265625,33.79443359375,33.79443359375,133,"[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT32x64x128_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA942_IU1_K1_LBSPPA512_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT2_1_MO40_NTn1_NTA4_NTB0_NTC0_NTD0_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW2_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA2_VWB1_WSGRA0_WSGRB0_WS64_WG16_16_1', 'stream': 0, 'count': 4, 'total_duration_us': np.float64(33.794), 'mean_duration_us': np.float64(8.4485), 'median_duration_us': np.float64(7.276), 'std_dev_duration_us': np.float64(2.1389002431156063), 'min_duration_us': np.float64(7.095), 'max_duration_us': np.float64(12.147)}]","[{'name': 'Cijk_Alik_Bljk_B_BS_BH_Bias_HA_S_SAV_UserArgs_MT32x64x128_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(8.45)}]",8.648639867641858,69.79055373738055 diff --git a/tests/traces/torch_compile_triton/trace_perf_report_csvs/SDPA_fwd.csv b/tests/traces/torch_compile_triton/trace_perf_report_csvs/SDPA_fwd.csv index bfc9efc1f..c09e96bdc 100644 --- a/tests/traces/torch_compile_triton/trace_perf_report_csvs/SDPA_fwd.csv +++ b/tests/traces/torch_compile_triton/trace_perf_report_csvs/SDPA_fwd.csv @@ -1,2 +1,3 @@ name,param: B,param: N_Q,param: H_Q,param: N_KV,param: H_KV,param: d_h_qk,param: d_h_v,param: dropout,param: causal,param: flash_impl,param: dtype_A_B,GFLOPS_first,Data Moved (MB)_first,FLOPS/Byte_first,TB/s_mean,TB/s_median,TB/s_std,TB/s_min,TB/s_max,TFLOPS/s_mean,TFLOPS/s_median,TFLOPS/s_std,TFLOPS/s_min,TFLOPS/s_max,process_name_first,process_label_first,thread_name_first,Compute Spec,kernel_details__summarize_kernel_stats,trunc_kernel_details,Input Dims_first,Input type_first,Input Strides_first,Concrete Inputs_first,Kernel Time (µs)_mean,Kernel Time (µs)_median,Kernel Time (µs)_std,Kernel Time (µs)_min,Kernel Time (µs)_max,Kernel Time (µs)_sum,name_count,UID_first aten::_scaled_dot_product_flash_attention,8,4096,16,4096,16,128,128,0.0,False,True,"('c10::BFloat16', 'c10::BFloat16')",1099.511627776,512.0,2048.0,0.14574986588715064,0.14574986588715064,,0.14574986588715064,0.14574986588715064,298.4957253368845,298.4957253368845,,298.4957253368845,298.4957253368845,python,CPU,thread 3182306 (python),matrix_bf16,"[{'name': 'attn_fwd', 'stream': 0, 'count': 1, 'total_duration_us': np.float64(3683.509), 'mean_duration_us': np.float64(3683.509), 'median_duration_us': np.float64(3683.509), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(3683.509), 'max_duration_us': np.float64(3683.509)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(3683.51)}]","[[8, 16, 4096, 128], [8, 16, 4096, 128], [8, 16, 4096, 128], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', 'Scalar', 'Scalar', 'Scalar', 'Scalar']","[[8388608, 128, 2048, 1], [8388608, 128, 2048, 1], [8388608, 128, 2048, 1], [], [], [], []]","['', '', '', '0.', 'False', 'False', '0.088388347648318433']",3683.5087890625,3683.5087890625,,3683.5087890625,3683.5087890625,3683.5087890625,1,9 +aten::_flash_attention_forward,8,4096,16,4096,16,128,128,0.0,False,True,nan,1099.511627776,512.0,2048.0,0.14574986588715064,0.14574986588715064,,0.14574986588715064,0.14574986588715064,298.4957253368845,298.4957253368845,,298.4957253368845,298.4957253368845,python,CPU,thread 3182306 (python),,"[{'name': 'attn_fwd', 'stream': 0, 'count': 1, 'total_duration_us': np.float64(3683.509), 'mean_duration_us': np.float64(3683.509), 'median_duration_us': np.float64(3683.509), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(3683.509), 'max_duration_us': np.float64(3683.509)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(3683.51)}]","[[8, 4096, 16, 128], [8, 4096, 16, 128], [8, 4096, 16, 128], [], [], [], [], [], [], [], [], [], [], [], []]","['c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '']","[[8388608, 2048, 128, 1], [8388608, 2048, 128, 1], [8388608, 2048, 128, 1], [], [], [], [], [], [], [], [], [], [], [], []]","['', '', '', '', '', '4096', '4096', '0.', 'False', 'False', '0.088388347648318433', '', '', '', '']",3683.5087890625,3683.5087890625,,3683.5087890625,3683.5087890625,3683.5087890625,1,16 diff --git a/tests/traces/torch_compile_triton/trace_perf_report_csvs/kernel_summary.csv b/tests/traces/torch_compile_triton/trace_perf_report_csvs/kernel_summary.csv index fdd6998dd..68a45d1f8 100644 --- a/tests/traces/torch_compile_triton/trace_perf_report_csvs/kernel_summary.csv +++ b/tests/traces/torch_compile_triton/trace_perf_report_csvs/kernel_summary.csv @@ -1,6 +1,6 @@ Parent op category,Parent cpu_op,Kernel name,Kernel stream,Kernel duration (µs)_sum,Kernel duration (µs)_count,Kernel duration (µs)_mean,Kernel duration (µs)_min,Kernel duration (µs)_max,Percent of kernels time (%),Percent of total time (%) GEMM,aten::mm,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMBSK_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_6_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO1_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0,3749.165,2,1874.5825,1775.035,1974.13,29.87615527936402,29.87612855168348 -other,aten::_flash_attention_forward,attn_fwd,0,3683.509,1,3683.509,3683.509,3683.509,29.3529590874061,29.352932827785132 +SDPA_fwd,aten::_flash_attention_forward,attn_fwd,0,3683.509,1,3683.509,3683.509,3683.509,29.3529590874061,29.352932827785132 GEMM,aten::mm,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMBSK_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_12_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG64_4_1,0,2246.691,4,561.67275,527.342,597.526,17.903316920100774,17.903300903510594 GEMM,aten::mm,Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMBSK_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_6_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1,0,2033.699,1,2033.699,2033.699,2033.699,16.206037108392756,16.20602261021591 triton,triton_poi_fused__unsafe_view_mul_silu_2,triton_poi_fused__unsafe_view_mul_silu_2,0,527.382,1,527.382,527.382,527.382,4.202574846276852,4.202571086586996 diff --git a/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary.csv b/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary.csv index 73a61d0ea..291b3dc61 100644 --- a/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary.csv +++ b/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary.csv @@ -1,6 +1,6 @@ name,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,total_direct_kernel_time_ms,total_subtree_kernel_time_ms,Count,Categories,Percentage (%),Cumulative Percentage (%) aten::mm,8029.55517578125,8029.55517578125,8.02955517578125,8.02955517578125,7,['GEMM'],63.9855107285316,63.9855107285316 -aten::_flash_attention_forward,3683.5087890625,3683.5087890625,3.6835087890625,3.6835087890625,1,['other'],29.35295741563505,93.33846814416665 +aten::_flash_attention_forward,3683.5087890625,3683.5087890625,3.6835087890625,3.6835087890625,1,['SDPA_fwd'],29.35295741563505,93.33846814416665 triton_poi_fused__unsafe_view_mul_silu_2,527.3818359375,527.3818359375,0.5273818359375,0.5273818359375,1,['triton'],4.202573540212125,97.54104168437878 triton_poi_fused__unsafe_view_add_3,140.84814453125,140.84814453125,0.14084814453125,0.14084814453125,1,['triton'],1.1223835275683505,98.66342521194713 triton_red_fused__unsafe_view_add_mean_mul_pow_rsqrt_1,90.65283203125,90.65283203125,0.09065283203125,0.09065283203125,1,['triton'],0.7223896753337838,99.38581488728092 diff --git a/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary_by_category.csv b/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary_by_category.csv index 4c484d8e9..a8971ff0a 100644 --- a/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary_by_category.csv +++ b/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_summary_by_category.csv @@ -1,4 +1,4 @@ op category,Count,total_direct_kernel_time_ms,Percentage (%),Cumulative Percentage (%) GEMM,7,8.02955517578125,63.9855107285316,63.9855107285316 -other,1,3.6835087890625,29.35295741563505,93.33846814416665 +SDPA_fwd,1,3.6835087890625,29.35295741563505,93.33846814416665 triton,4,0.83595703125,6.661531855833344,100.0 diff --git a/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_unique_args.csv b/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_unique_args.csv index b23239a6e..5d8420c75 100644 --- a/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_unique_args.csv +++ b/tests/traces/torch_compile_triton/trace_perf_report_csvs/ops_unique_args.csv @@ -1,6 +1,6 @@ name,op category,process_name,process_label,thread_name,Input Dims,Input type,Input Strides,Concrete Inputs,operation_count,total_direct_kernel_time_mean,total_subtree_kernel_time_mean,total_direct_kernel_time_median,total_subtree_kernel_time_median,total_direct_kernel_time_std,total_subtree_kernel_time_std,total_direct_kernel_time_min,total_subtree_kernel_time_min,total_direct_kernel_time_max,total_subtree_kernel_time_max,total_direct_kernel_time_sum,total_subtree_kernel_time_sum,ex_UID,kernel_details_summary,trunc_kernel_details,Percentage (%),Cumulative Percentage (%) aten::mm,GEMM,python,CPU,thread 3182306 (python),"((32768, 2048), (2048, 8192), (32768, 8192))","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16')","((2048, 1), (1, 2048), (8192, 1))","('', '', '')",2,1874.58251953125,1874.58251953125,1874.58251953125,1874.58251953125,140.7812312508252,140.7812312508252,1775.03515625,1775.03515625,1974.1298828125,1974.1298828125,3749.1650390625,3749.1650390625,37,"[{'name': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMBSK_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_6_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO1_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 2, 'total_duration_us': np.float64(3749.165), 'mean_duration_us': np.float64(1874.5825), 'median_duration_us': np.float64(1874.5825), 'std_dev_duration_us': np.float64(99.5475), 'min_duration_us': np.float64(1775.035), 'max_duration_us': np.float64(1974.13)}]","[{'name': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(1874.58)}]",29.876155599943115,29.876155599943115 -aten::_flash_attention_forward,other,python,CPU,thread 3182306 (python),"((8, 4096, 16, 128), (8, 4096, 16, 128), (8, 4096, 16, 128), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((8388608, 2048, 128, 1), (8388608, 2048, 128, 1), (8388608, 2048, 128, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '4096', '4096', '0.', 'False', 'False', '0.088388347648318433', '', '', '', '')",1,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,,,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,16,"[{'name': 'attn_fwd', 'stream': 0, 'count': 1, 'total_duration_us': np.float64(3683.509), 'mean_duration_us': np.float64(3683.509), 'median_duration_us': np.float64(3683.509), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(3683.509), 'max_duration_us': np.float64(3683.509)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(3683.51)}]",29.35295741563505,59.229113015578164 +aten::_flash_attention_forward,SDPA_fwd,python,CPU,thread 3182306 (python),"((8, 4096, 16, 128), (8, 4096, 16, 128), (8, 4096, 16, 128), (), (), (), (), (), (), (), (), (), (), (), ())","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16', '', '', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', 'Scalar', '', '', '', '')","((8388608, 2048, 128, 1), (8388608, 2048, 128, 1), (8388608, 2048, 128, 1), (), (), (), (), (), (), (), (), (), (), (), ())","('', '', '', '', '', '4096', '4096', '0.', 'False', 'False', '0.088388347648318433', '', '', '', '')",1,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,,,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,3683.5087890625,16,"[{'name': 'attn_fwd', 'stream': 0, 'count': 1, 'total_duration_us': np.float64(3683.509), 'mean_duration_us': np.float64(3683.509), 'median_duration_us': np.float64(3683.509), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(3683.509), 'max_duration_us': np.float64(3683.509)}]","[{'name': 'attn_fwd', 'stream': 0, 'mean_duration_us': np.float64(3683.51)}]",29.35295741563505,59.229113015578164 aten::mm,GEMM,python,CPU,thread 3182306 (python),"((32768, 2048), (2048, 2048), (32768, 2048))","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16')","((2048, 1), (1, 2048), (2048, 1))","('', '', '')",4,561.6727294921875,561.6727294921875,560.91162109375,560.91162109375,28.81344828937993,28.81344828937993,527.341796875,527.341796875,597.52587890625,597.52587890625,2246.69091796875,2246.69091796875,6,"[{'name': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMBSK_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA512_LBSPPB512_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT4_12_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW4_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA4_VWB4_WSGRA0_WSGRB0_WS64_WG64_4_1', 'stream': 0, 'count': 4, 'total_duration_us': np.float64(2246.691), 'mean_duration_us': np.float64(561.67275), 'median_duration_us': np.float64(560.9114999999999), 'std_dev_duration_us': np.float64(24.95315869158651), 'min_duration_us': np.float64(527.342), 'max_duration_us': np.float64(597.526)}]","[{'name': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(561.67)}]",17.90331627198726,77.13242928756543 aten::mm,GEMM,python,CPU,thread 3182306 (python),"((32768, 8192), (8192, 2048), (32768, 2048))","('c10::BFloat16', 'c10::BFloat16', 'c10::BFloat16')","((8192, 1), (1, 8192), (2048, 1))","('', '', '')",1,2033.69921875,2033.69921875,2033.69921875,2033.69921875,,,2033.69921875,2033.69921875,2033.69921875,2033.69921875,2033.69921875,2033.69921875,40,"[{'name': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTLA0_DTLB0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMBSK_GLS0_ISA942_IU1_K1_LDSTI0_LBSPPA1024_LBSPPB256_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT8_6_MO40_NTn1_NTA0_NTB0_NTC0_NTD4_NTM0_NEPBS16_NLCA1_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW8_SK0_SKFTR0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA8_VWB2_WSGRA0_WSGRB0_WS64_WG32_8_1', 'stream': 0, 'count': 1, 'total_duration_us': np.float64(2033.699), 'mean_duration_us': np.float64(2033.699), 'median_duration_us': np.float64(2033.699), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(2033.699), 'max_duration_us': np.float64(2033.699)}]","[{'name': 'Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT256x192x64_MI16x1...', 'stream': 0, 'mean_duration_us': np.float64(2033.7)}]",16.206038856601232,93.33846814416665 triton_poi_fused__unsafe_view_mul_silu_2,triton,python,CPU,thread 3182306 (python),"((8, 4096, 8192), (32768, 8192), ())","('c10::BFloat16', 'c10::BFloat16', 'Scalar')","((33554432, 8192, 1), (8192, 1), ())","('', '', '268435456')",1,527.3818359375,527.3818359375,527.3818359375,527.3818359375,,,527.3818359375,527.3818359375,527.3818359375,527.3818359375,527.3818359375,527.3818359375,39,"[{'name': 'triton_poi_fused__unsafe_view_mul_silu_2', 'stream': 0, 'count': 1, 'total_duration_us': np.float64(527.382), 'mean_duration_us': np.float64(527.382), 'median_duration_us': np.float64(527.382), 'std_dev_duration_us': np.float64(0.0), 'min_duration_us': np.float64(527.382), 'max_duration_us': np.float64(527.382)}]","[{'name': 'triton_poi_fused__unsafe_view_mul_silu_2', 'stream': 0, 'mean_duration_us': np.float64(527.38)}]",4.202573540212126,97.54104168437878