diff --git a/tests/cpu/test_autotune_session_split.py b/tests/cpu/test_autotune_session_split.py new file mode 100644 index 0000000..7ba4a66 --- /dev/null +++ b/tests/cpu/test_autotune_session_split.py @@ -0,0 +1,647 @@ +# Copyright (c) Meta Platforms, Inc. and affiliates. + +"""Autotune sessions must split by key, not just by call site. + +Two ``Autotuner.run()`` invocations at the same call site with different keys +(e.g. ``M=N=16`` then ``M=N=32`` in a loop) share one call-site session id. +The parser must still emit one ``autotune_analysis`` event per key round, so +per-config arguments stay single-valued in each session's varies table. +""" + +import json +import os +import tempfile +import unittest + +from tritonparse.parse.event_diff import ( + _config_value_matches, + _extract_scalar_value, + _parse_best_config, +) +from tritonparse.parse.trace_processor import parse_single_rank + + +def _stack(boundary: str) -> list[dict]: + # Identical user frame for every event: both key rounds run at the same + # call site, so they share one coarse (call-site) session id. + return [ + {"filename": "/user/loop.py", "name": "run_kernel", "line": 42}, + {"filename": "triton/runtime/autotuner.py", "name": boundary, "line": 100}, + ] + + +def _compilation(kernel_hash: str) -> dict: + return { + "event_type": "compilation", + "pid": 1000, + "timestamp": "2026-09-02T00:00:00", + "stack": _stack("_bench"), + "payload": { + "metadata": { + "hash": kernel_hash, + "name": "matmul_kernel", + "num_warps": 1, + "num_stages": 1, + "num_ctas": 1, + }, + "pt_info": {"frame_id": 0, "frame_compile_id": 0}, + "file_content": {}, + "file_path": {}, + }, + } + + +def _extracted_args(m: int, block_m: int, data_ptr: str) -> dict: + return { + "a": { + "type": "tensor", + "shape": [m, 16], + "dtype": "torch.float16", + "stride": [16, 1], + "data_ptr": data_ptr, + }, + "M": {"type": "int", "value": m}, + "N": {"type": "int", "value": m}, + "BLOCK_SIZE_M": {"type": "int", "value": block_m}, + } + + +def _launch( + kernel_hash: str, + m: int, + block_m: int, + boundary: str, + data_ptr: str, +) -> dict: + return { + "event_type": "launch", + "pid": 1000, + "timestamp": "2026-09-02T00:00:01", + "name": "matmul_kernel", + "stack": _stack(boundary), + "compilation_metadata": { + "hash": kernel_hash, + "name": "matmul_kernel", + "num_warps": 1, + "num_stages": 1, + "num_ctas": 1, + }, + "extracted_args": _extracted_args(m, block_m, data_ptr), + } + + +def _autotune_event(m: int, block_m: int) -> dict: + return { + "event_type": "autotune", + "kernel_name": "matmul_kernel", + "stack": _stack("run"), + "best_config": ( + f"BLOCK_SIZE_M: {block_m}, BLOCK_SIZE_N: 16, " + "num_warps: 1, num_stages: 1, num_ctas: 1" + ), + "configs_timings": {"cfg": 1.0}, + "duration": 0.5, + "cache_hit": False, + "cache_key": "cache-key", + "autotune_key": f"({m}, {m}, 'torch.float16')", + } + + +def _write_trace(path: str, events: list[dict]) -> None: + with open(path, "w") as output: + for event in events: + output.write(json.dumps(event) + "\n") + + +def _read_output_events(output_dir: str) -> list[dict]: + events = [] + for filename in sorted(os.listdir(output_dir)): + if filename.endswith(".ndjson"): + with open(os.path.join(output_dir, filename), "r") as source: + events.extend(json.loads(line) for line in source if line.strip()) + return events + + +def _two_key_trace() -> list[dict]: + """Two key rounds (M=16, M=32) at one call site, two configs each.""" + return [ + _compilation("hash_a"), + _compilation("hash_b"), + # Round 1: M = N = 16, winner is hash_a (BLOCK_SIZE_M 16). + _launch("hash_a", 16, 16, "_bench", "0x1000"), + _launch("hash_b", 16, 32, "_bench", "0x2000"), + _launch("hash_a", 16, 16, "run", "0x3000"), + _autotune_event(16, 16), + # Round 2: M = N = 32, winner is hash_b (BLOCK_SIZE_M 32). + _launch("hash_a", 32, 16, "_bench", "0x4000"), + _launch("hash_b", 32, 32, "_bench", "0x5000"), + _launch("hash_b", 32, 32, "run", "0x6000"), + _autotune_event(32, 32), + ] + + +class AutotuneSessionSplitTest(unittest.TestCase): + def setUp(self) -> None: + self.temporary_directory = tempfile.TemporaryDirectory() + self.input_path = os.path.join( + self.temporary_directory.name, + "dedicated_log_triton_trace_user_pid_1000_.ndjson", + ) + self.output_dir = os.path.join(self.temporary_directory.name, "output") + os.makedirs(self.output_dir) + + def tearDown(self) -> None: + self.temporary_directory.cleanup() + + def test_two_keys_share_call_site_but_split_into_two_sessions(self) -> None: + _write_trace(self.input_path, _two_key_trace()) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 2) + by_winner = {a["winner_compilation_hash"]: a for a in analyses} + self.assertEqual(set(by_winner), {"hash_a", "hash_b"}) + + # Each key round keeps single-valued args: M/N are constant + # within the round (scalars in the sames section), only the true + # config param varies, and no cell is a multi-valued + # {unique_count, values} distribution (which the viewer renders + # as an ellipsis). + for winner_hash, analysis in by_winner.items(): + configs = analysis["autotune_args_summary"]["autotune_configs"] + varies = configs["varies"] + sames = configs["sames"] + expected_m = 16 if winner_hash == "hash_a" else 32 + self.assertEqual(sames["M"], {"type": "int", "value": expected_m}) + self.assertEqual(sames["N"], {"type": "int", "value": expected_m}) + self.assertEqual(set(varies), {"BLOCK_SIZE_M"}) + self.assertEqual( + varies["BLOCK_SIZE_M"]["hash_a"], {"type": "int", "value": 16} + ) + self.assertEqual( + varies["BLOCK_SIZE_M"]["hash_b"], {"type": "int", "value": 32} + ) + self.assertNotIn("unique_count", json.dumps(configs)) + + # Winners and occurrence ids stay with their own key round. + self.assertEqual(by_winner["hash_a"]["launch_occurrence_ids"]["winner"], [4]) + self.assertEqual(by_winner["hash_b"]["launch_occurrence_ids"]["winner"], [7]) + self.assertEqual( + sorted(by_winner["hash_a"]["launch_occurrence_ids"]["benchmark"]), + [2, 3], + ) + self.assertEqual( + sorted(by_winner["hash_b"]["launch_occurrence_ids"]["benchmark"]), + [5, 6], + ) + + # Each session keeps its own autotune result, matched by best config. + self.assertIn( + "BLOCK_SIZE_M: 16", by_winner["hash_a"]["autotune_result"]["best_config"] + ) + self.assertIn( + "BLOCK_SIZE_M: 32", by_winner["hash_b"]["autotune_result"]["best_config"] + ) + + # The two sub-sessions are distinguishable by their full key + # signatures (never truncated: truncation could collide and drop + # a key round). + session_ids = {a["session_id"] for a in analyses} + self.assertEqual(len(session_ids), 2) + for session_id in session_ids: + coarse, _, sig = session_id.partition(":") + self.assertTrue(coarse) + self.assertEqual(len(sig), 16) + + def test_two_callbacks_without_args_share_one_partition(self) -> None: + # Degraded but valid trace shape: two key rounds (distinct + # listener callbacks) whose launches carry no extracted_args. + # The launches are indistinguishable, so no key attribution is + # possible: expect one analysis with the latest result attached. + events = _two_key_trace() + for event in events: + if event.get("event_type") == "launch": + del event["extracted_args"] + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + self.assertNotIn(":", analyses[0]["session_id"]) + self.assertEqual(analyses[0]["winner_compilation_hash"], "hash_b") + self.assertIn("BLOCK_SIZE_M: 32", analyses[0]["autotune_result"]["best_config"]) + + def test_split_preserves_compilation_without_traced_launches(self) -> None: + # A candidate whose benchmark launches were never traced cannot be + # attributed to a key round, but the split must not silently drop + # it: every partition keeps it alongside its referenced configs. + events = _two_key_trace() + events.insert(2, _compilation("hash_c")) + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 2) + for analysis in analyses: + self.assertEqual( + analysis["compilation_analysis"]["compilation_hashes"], + ["hash_a", "hash_b", "hash_c"], + ) + # hash_c has no launches, so it contributes no per-config args. + varies = analysis["autotune_args_summary"]["autotune_configs"]["varies"] + self.assertNotIn("hash_c", varies["BLOCK_SIZE_M"]) + by_winner = {a["winner_compilation_hash"]: a for a in analyses} + self.assertEqual(set(by_winner), {"hash_a", "hash_b"}) + + def test_note_only_launches_are_ignored_for_splitting(self) -> None: + # Launches captured inside CUDA graphs record a _note marker + # instead of real args. They hold no key information and must + # neither split their own round nor pool unrelated rounds; the + # same rounds' warmup launches outside capture still split + # normally, and note occurrences stay out of the partitions. + events = _two_key_trace() + note_a = _launch("hash_a", 16, 16, "_bench", "0x1000") + note_a["extracted_args"] = {"_note": "argument extraction skipped"} + note_b = _launch("hash_b", 32, 32, "_bench", "0x5000") + note_b["extracted_args"] = {"_note": "argument extraction skipped"} + events.insert(3, note_a) + events.insert(8, note_b) + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 2) + by_winner = {a["winner_compilation_hash"]: a for a in analyses} + self.assertEqual(set(by_winner), {"hash_a", "hash_b"}) + # Warmups only; note occurrences (3 and 7) are excluded. + self.assertEqual( + sorted(by_winner["hash_a"]["launch_occurrence_ids"]["benchmark"]), + [2, 4], + ) + self.assertEqual( + sorted(by_winner["hash_b"]["launch_occurrence_ids"]["benchmark"]), + [6, 8], + ) + self.assertNotIn( + "unique_count", + json.dumps( + by_winner["hash_a"]["autotune_args_summary"]["autotune_configs"] + ), + ) + + def test_single_partition_keeps_note_occurrences(self) -> None: + # One key round mixed with graph-captured launches: no + # attribution choice exists, so note occurrences join the + # emission (benchmark ids [2, 3, 4], winner ids [5, 6]) while + # their _note payload stays out of the varies table. + events = _two_key_trace()[:6] + note_bench = _launch("hash_a", 16, 16, "_bench", "0x1000") + note_bench["extracted_args"] = {"_note": "captured"} + note_winner = _launch("hash_a", 16, 16, "run", "0x3000") + note_winner["extracted_args"] = {"_note": "captured"} + events.insert(2, note_bench) + events.insert(6, note_winner) + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + self.assertNotIn(":", analyses[0]["session_id"]) + self.assertEqual( + sorted(analyses[0]["launch_occurrence_ids"]["benchmark"]), + [2, 3, 4], + ) + self.assertEqual(sorted(analyses[0]["launch_occurrence_ids"]["winner"]), [5, 6]) + self.assertFalse(analyses[0]["cache_usage"]) + self.assertNotIn( + "_note", + json.dumps(analyses[0]["autotune_args_summary"]["autotune_configs"]), + ) + + def test_all_note_session_falls_back_to_coarse_occurrences(self) -> None: + # Every launch captured inside CUDA graphs: no key derivation is + # possible, but the emission must keep the real launches via the + # coarse occurrence ids instead of reporting cache_usage. + events = _two_key_trace()[:6] + for event in events: + if event.get("event_type") == "launch": + event["extracted_args"] = {"_note": "captured"} + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + self.assertEqual( + sorted(analyses[0]["launch_occurrence_ids"]["benchmark"]), [2, 3] + ) + self.assertEqual(sorted(analyses[0]["launch_occurrence_ids"]["winner"]), [4]) + self.assertFalse(analyses[0]["cache_usage"]) + self.assertEqual(analyses[0]["winner_compilation_hash"], "hash_a") + self.assertIn("autotune_result", analyses[0]) + + def test_optional_scalar_config_stays_whole(self) -> None: + # An optional config arg recorded as None under one compilation + # ({"type": "NoneType", "repr": "None"}) and as an int under the + # other ({"type": "int", "value": 16}) must still be recognized + # as one config param: a nested path missing under a hash counts + # as a distinct value, so the whole wrapper is masked instead of + # splitting the round by config (each half would then fall below + # the emission threshold and drop the entire session). + events = _two_key_trace()[:6] + for event in events: + if event.get("event_type") == "launch": + if event["compilation_metadata"]["hash"] == "hash_a": + event["extracted_args"]["OPTIONAL_TILE"] = { + "type": "NoneType", + "repr": "None", + } + else: + event["extracted_args"]["OPTIONAL_TILE"] = { + "type": "int", + "value": 16, + } + elif event.get("event_type") == "autotune": + event["best_config"] += ", OPTIONAL_TILE: None" + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + self.assertNotIn(":", analyses[0]["session_id"]) + self.assertEqual(analyses[0]["winner_compilation_hash"], "hash_a") + self.assertIn( + "OPTIONAL_TILE: None", + analyses[0]["autotune_result"]["best_config"], + ) + + def test_all_note_session_with_external_compilations(self) -> None: + # Kernels compiled outside autotuning (no autotuner frames on the + # compilation stack) leave the session with no compilation events. + # The all-note fallback preserves the coarse launch groups, so + # output-file lookup resolves through the launches' compilation + # hashes and the session still emits instead of hitting the + # neither-compilations-nor-launches guard. + events = _two_key_trace()[:6] + for event in events: + if event.get("event_type") == "compilation": + event["stack"] = [ + {"filename": "/user/loop.py", "name": "run_kernel", "line": 1} + ] + elif event.get("event_type") == "launch": + event["extracted_args"] = {"_note": "captured"} + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + self.assertEqual( + sorted(analyses[0]["launch_occurrence_ids"]["benchmark"]), [2, 3] + ) + self.assertEqual(sorted(analyses[0]["launch_occurrence_ids"]["winner"]), [4]) + self.assertFalse(analyses[0]["cache_usage"]) + self.assertEqual(analyses[0]["winner_compilation_hash"], "hash_a") + self.assertIn("autotune_result", analyses[0]) + + def test_mixed_descriptor_masks_only_config_paths(self) -> None: + # A descriptor mixing key fields (shape, one per round) with + # config fields (block_shape, one per compilation) must split on + # the key fields only. Masking the whole argument would merge the + # rounds; keeping it whole would over-split on block_shape (four + # partitions); per-path masking yields exactly the two rounds. + def _descriptor(shape, block_shape, data_ptr): + return { + "type": "TensorDescriptor", + "base": { + "type": "tensor", + "shape": [16, 16], + "dtype": "torch.float16", + "stride": [16, 1], + "data_ptr": data_ptr, + }, + "shape": shape, + "strides": [16, 1], + "block_shape": block_shape, + "padding": "zero", + } + + def _desc_launch(kernel_hash, shape, block_shape, boundary, data_ptr): + launch = _launch( + kernel_hash, + 16, + 16 if kernel_hash == "hash_a" else 32, + boundary, + data_ptr, + ) + launch["extracted_args"]["desc"] = _descriptor(shape, block_shape, data_ptr) + return launch + + events = [ + _compilation("hash_a"), + _compilation("hash_b"), + # Round 1: shape [16, 16], winner hash_a. + _desc_launch("hash_a", [16, 16], [16, 16], "_bench", "0x1000"), + _desc_launch("hash_b", [16, 16], [32, 16], "_bench", "0x2000"), + _desc_launch("hash_a", [16, 16], [16, 16], "run", "0x3000"), + _autotune_event(16, 16), + # Round 2: shape [32, 32], winner hash_b. + _desc_launch("hash_a", [32, 32], [16, 16], "_bench", "0x4000"), + _desc_launch("hash_b", [32, 32], [32, 16], "_bench", "0x5000"), + _desc_launch("hash_b", [32, 32], [32, 16], "run", "0x6000"), + _autotune_event(16, 32), + ] + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 2) + by_winner = {a["winner_compilation_hash"]: a for a in analyses} + self.assertEqual(set(by_winner), {"hash_a", "hash_b"}) + self.assertIn( + "BLOCK_SIZE_M: 16", by_winner["hash_a"]["autotune_result"]["best_config"] + ) + self.assertIn( + "BLOCK_SIZE_M: 32", by_winner["hash_b"]["autotune_result"]["best_config"] + ) + for analysis in analyses: + _, _, sig = analysis["session_id"].partition(":") + self.assertEqual(len(sig), 16) + + def test_descriptor_args_stable_across_reallocation(self) -> None: + # TMA descriptors nest backing-tensor data_ptrs. Reallocating the + # backing tensor must not split a round, and config-shaped fields + # (block_shape rewritten per config) must still be recognized as + # config params so the round stays whole. + def _descriptor(block_shape, data_ptr): + return { + "type": "TensorDescriptor", + "base": { + "type": "tensor", + "shape": [16, 16], + "dtype": "torch.float16", + "stride": [16, 1], + "data_ptr": data_ptr, + }, + "shape": [16, 16], + "strides": [16, 1], + "block_shape": block_shape, + "padding": "zero", + } + + def _desc_launch(kernel_hash, block_shape, data_ptr, boundary): + launch = _launch(kernel_hash, 16, 16, boundary, data_ptr) + launch["extracted_args"]["desc"] = _descriptor(block_shape, data_ptr) + return launch + + events = [ + _compilation("hash_a"), + _compilation("hash_b"), + _desc_launch("hash_a", [16, 16], "0x1000", "_bench"), + _desc_launch("hash_b", [32, 16], "0x1000", "_bench"), + # Winner rerun after the backing tensor was reallocated. + _desc_launch("hash_a", [16, 16], "0x9999", "run"), + _autotune_event(16, 16), + ] + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + self.assertNotIn(":", analyses[0]["session_id"]) + self.assertEqual(analyses[0]["winner_compilation_hash"], "hash_a") + self.assertIn("autotune_result", analyses[0]) + + def test_warp_specialization_matches_on_warps_base(self) -> None: + # Under warp specialization the launch metadata carries the + # expanded num_warps while best_config prints the requested one; + # matching must use the compilation payload's num_warps_base. + events = _two_key_trace() + for event in events: + if event.get("event_type") == "launch": + event["compilation_metadata"]["num_warps"] = 8 + elif event.get("event_type") == "compilation": + event["payload"]["metadata"]["num_warps"] = 8 + event["payload"]["metadata"]["num_warps_base"] = 4 + elif event.get("event_type") == "autotune": + event["best_config"] = event["best_config"].replace( + "num_warps: 1", "num_warps: 4" + ) + _write_trace(self.input_path, events) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 2) + by_winner = {a["winner_compilation_hash"]: a for a in analyses} + self.assertIn( + "BLOCK_SIZE_M: 16", by_winner["hash_a"]["autotune_result"]["best_config"] + ) + self.assertIn( + "BLOCK_SIZE_M: 32", by_winner["hash_b"]["autotune_result"]["best_config"] + ) + + def test_single_key_session_keeps_plain_session_id(self) -> None: + _write_trace(self.input_path, _two_key_trace()[:6]) + parse_single_rank([self.input_path], self.output_dir) + + analyses = [ + e + for e in _read_output_events(self.output_dir) + if e.get("event_type") == "autotune_analysis" + ] + self.assertEqual(len(analyses), 1) + # No split happened, so the session id keeps its legacy plain form. + self.assertNotIn(":", analyses[0]["session_id"]) + self.assertEqual(analyses[0]["winner_compilation_hash"], "hash_a") + self.assertIn("autotune_result", analyses[0]) + + +class BestConfigParsingTest(unittest.TestCase): + def test_tuple_value_with_commas_survives(self) -> None: + parsed = _parse_best_config("BLOCK: (16, 32), BLOCK_SIZE_M: 16, num_warps: 1") + self.assertEqual( + parsed, + {"BLOCK": "(16, 32)", "BLOCK_SIZE_M": "16", "num_warps": "1"}, + ) + + def test_non_string_config_parses_to_empty(self) -> None: + self.assertEqual(_parse_best_config(None), {}) + self.assertEqual(_parse_best_config({"BLOCK_SIZE_M": 16}), {}) + + def test_config_value_matching(self) -> None: + # Exact ints, including beyond float precision and hex spellings. + self.assertTrue(_config_value_matches("16", 16)) + self.assertTrue(_config_value_matches("0x10", 16)) + self.assertTrue(_config_value_matches(str(2**60 + 1), 2**60 + 1)) + self.assertFalse(_config_value_matches("17", 16)) + # None spellings. + self.assertTrue(_config_value_matches("None", None)) + self.assertFalse(_config_value_matches("16", None)) + # Bools only accept true/false/1/0 spellings, never other numbers. + self.assertTrue(_config_value_matches("True", True)) + self.assertTrue(_config_value_matches("1", True)) + self.assertTrue(_config_value_matches("False", False)) + self.assertTrue(_config_value_matches("0", False)) + self.assertFalse(_config_value_matches("True", 1)) + self.assertFalse(_config_value_matches("2", True)) + self.assertFalse(_config_value_matches("-1", True)) + self.assertFalse(_config_value_matches("yes", True)) + self.assertFalse(_config_value_matches("True", False)) + # Floats and plain strings fall through. + self.assertTrue(_config_value_matches("1.5", 1.5)) + self.assertTrue(_config_value_matches("relu", "relu")) + self.assertFalse(_config_value_matches("relu", "none")) + # Sequences: best_config prints tuples while JSON round-trips + # turn them into lists. + self.assertTrue(_config_value_matches("(2, 1, 1)", [2, 1, 1])) + self.assertTrue(_config_value_matches("(2, 1, 1)", (2, 1, 1))) + self.assertFalse(_config_value_matches("(2, 1)", [2, 1, 1])) + self.assertFalse(_config_value_matches("None", [1])) + # "repr" payloads unwrap for values without structured form. + self.assertEqual( + _extract_scalar_value({"type": "NoneType", "repr": "None"}), "None" + ) + self.assertTrue(_config_value_matches("None", "None")) + self.assertTrue(_config_value_matches("(16, 32)", "(16, 32)")) diff --git a/tests/cpu/test_kernel_query.py b/tests/cpu/test_kernel_query.py index caecc41..27120d7 100644 --- a/tests/cpu/test_kernel_query.py +++ b/tests/cpu/test_kernel_query.py @@ -86,7 +86,7 @@ def test_list_kernels_multiple(self): # half-written fixture; refresh it with `make regen-examples-install`. kernel_dict = {k.name: k for k in result} self.assertEqual(kernel_dict["fused_op_kernel"].total_launches, 4) - self.assertEqual(kernel_dict["matmul_kernel"].total_launches, 1078) + self.assertEqual(kernel_dict["matmul_kernel"].total_launches, 1041) def test_find_launch_index_valid(self): """Test finding valid kernel name and launch_id.""" diff --git a/tests/example_output/logs/dedicated_log_triton_trace_findhao_.ndjson b/tests/example_output/logs/dedicated_log_triton_trace_findhao_.ndjson index 68e1e4f..977990c 100644 --- a/tests/example_output/logs/dedicated_log_triton_trace_findhao_.ndjson +++ b/tests/example_output/logs/dedicated_log_triton_trace_findhao_.ndjson @@ -1,24 +1,26 @@ -{"event_type":"compilation","pid":1724463,"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":149,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":739,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":884,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[["tt.divisibility",16]],"(5,)":[["tt.divisibility",16]],"(6,)":[["tt.divisibility",16]],"(7,)":[],"(8,)":[["tt.divisibility",16]],"(9,)":[],"(10,)":[["tt.divisibility",16]],"(11,)":[],"(12,)":[],"(13,)":[],"(14,)":[],"(15,)":[],"(16,)":[]},"src_cache_key":"528f8911847b1f39f462157322ec9a674a5565b9b74d621d5ff09625eb710f38","src_constants":{"(7,)":1,"(9,)":1,"(11,)":1,"(12,)":16,"(13,)":16,"(14,)":16,"(15,)":1,"(16,)":null},"times":{"ir_initialization":377646,"lowering_stages":[],"store_results":0}},"file_path":{"matmul_kernel.source":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.source","matmul_kernel.ttir":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.ttir","matmul_kernel.ttgir":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.ttgir","matmul_kernel.llir":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.llir","matmul_kernel.ptx":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.ptx","matmul_kernel.cubin":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.cubin","matmul_kernel.json":"/home/user/.triton/cache/OUNNYLK5WWI2ZVOWQVDW2KQKYTYWY4C2H6C3O4KG3NCTLTMRRCFQ/matmul_kernel.json"},"file_content":{"matmul_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:0)\n#loc55 = loc(\"a\"(#loc))\n#loc56 = loc(\"b\"(#loc))\n#loc57 = loc(\"c\"(#loc))\n#loc58 = loc(\"M\"(#loc))\n#loc59 = loc(\"N\"(#loc))\n#loc60 = loc(\"K\"(#loc))\n#loc61 = loc(\"stride_am\"(#loc))\n#loc62 = loc(\"stride_bk\"(#loc))\n#loc63 = loc(\"stride_cm\"(#loc))\nmodule {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %cst = arith.constant dense<0.000000e+00> : tensor<16x16xf16> loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc2)\n %cst_0 = arith.constant dense<16> : tensor<16x16xi32> loc(#loc1)\n %cst_1 = arith.constant dense<0.000000e+00> : tensor<16x16xf32> loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc64)\n %num_pid_m = arith.addi %M, %c15_i32 : i32 loc(#loc113)\n %num_pid_m_2 = arith.divsi %num_pid_m, %c16_i32 : i32 loc(#loc114)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc115)\n %num_pid_n_3 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc116)\n %group_id = arith.divsi %pid, %num_pid_n_3 : i32 loc(#loc67)\n %group_size = arith.subi %num_pid_m_2, %group_id : i32 loc(#loc68)\n %group_size_4 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc69)\n %pid_m = arith.remsi %pid, %group_size_4 : i32 loc(#loc70)\n %pid_m_5 = arith.addi %group_id, %pid_m : i32 loc(#loc71)\n %pid_n = arith.remsi %pid, %num_pid_n_3 : i32 loc(#loc72)\n %pid_n_6 = arith.divsi %pid_n, %group_size_4 : i32 loc(#loc73)\n %offs_am = arith.muli %pid_m_5, %c16_i32 : i32 loc(#loc74)\n %offs_am_7 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc75)\n %offs_am_8 = tt.splat %offs_am : i32 -> tensor<16xi32> loc(#loc76)\n %offs_am_9 = arith.addi %offs_am_8, %offs_am_7 : tensor<16xi32> loc(#loc76)\n %offs_am_10 = tt.splat %M : i32 -> tensor<16xi32> loc(#loc77)\n %offs_am_11 = arith.remsi %offs_am_9, %offs_am_10 : tensor<16xi32> loc(#loc77)\n %offs_bn = arith.muli %pid_n_6, %c16_i32 : i32 loc(#loc78)\n %offs_bn_12 = tt.splat %offs_bn : i32 -> tensor<16xi32> loc(#loc79)\n %offs_bn_13 = arith.addi %offs_bn_12, %offs_am_7 : tensor<16xi32> loc(#loc79)\n %offs_bn_14 = tt.splat %N : i32 -> tensor<16xi32> loc(#loc80)\n %offs_bn_15 = arith.remsi %offs_bn_13, %offs_bn_14 : tensor<16xi32> loc(#loc80)\n %a_ptrs = tt.expand_dims %offs_am_11 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc81)\n %a_ptrs_16 = tt.splat %stride_am : i32 -> tensor<16x1xi32> loc(#loc82)\n %a_ptrs_17 = arith.muli %a_ptrs, %a_ptrs_16 : tensor<16x1xi32> loc(#loc82)\n %a_ptrs_18 = tt.expand_dims %offs_am_7 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc83)\n %a_ptrs_19 = tt.broadcast %a_ptrs_17 : tensor<16x1xi32> -> tensor<16x16xi32> loc(#loc84)\n %a_ptrs_20 = tt.broadcast %a_ptrs_18 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc84)\n %a_ptrs_21 = arith.addi %a_ptrs_19, %a_ptrs_20 : tensor<16x16xi32> loc(#loc84)\n %a_ptrs_22 = tt.splat %a : !tt.ptr -> tensor<16x16x!tt.ptr> loc(#loc85)\n %a_ptrs_23 = tt.addptr %a_ptrs_22, %a_ptrs_21 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc85)\n %b_ptrs = tt.expand_dims %offs_am_7 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc86)\n %b_ptrs_24 = tt.splat %stride_bk : i32 -> tensor<16x1xi32> loc(#loc87)\n %b_ptrs_25 = arith.muli %b_ptrs, %b_ptrs_24 : tensor<16x1xi32> loc(#loc87)\n %b_ptrs_26 = tt.expand_dims %offs_bn_15 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc88)\n %b_ptrs_27 = tt.broadcast %b_ptrs_25 : tensor<16x1xi32> -> tensor<16x16xi32> loc(#loc89)\n %b_ptrs_28 = tt.broadcast %b_ptrs_26 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc89)\n %b_ptrs_29 = arith.addi %b_ptrs_27, %b_ptrs_28 : tensor<16x16xi32> loc(#loc89)\n %b_ptrs_30 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr> loc(#loc90)\n %b_ptrs_31 = tt.addptr %b_ptrs_30, %b_ptrs_29 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc90)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc91)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc92)\n %accumulator:3 = scf.for %k = %c0_i32 to %1 step %c1_i32 iter_args(%a_ptrs_46 = %a_ptrs_23, %b_ptrs_47 = %b_ptrs_31, %accumulator_48 = %cst_1) -> (tensor<16x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<16x16xf32>) : i32 {\n %a_block = arith.muli %k, %c16_i32 : i32 loc(#loc94)\n %a_block_49 = arith.subi %K, %a_block : i32 loc(#loc95)\n %a_block_50 = tt.splat %a_block_49 : i32 -> tensor<1x16xi32> loc(#loc96)\n %a_block_51 = arith.cmpi slt, %a_ptrs_18, %a_block_50 : tensor<1x16xi32> loc(#loc96)\n %a_block_52 = tt.broadcast %a_block_51 : tensor<1x16xi1> -> tensor<16x16xi1> loc(#loc97)\n %a_block_53 = tt.load %a_ptrs_46, %a_block_52, %cst : tensor<16x16x!tt.ptr> loc(#loc97)\n %b_block = tt.splat %a_block_49 : i32 -> tensor<16x1xi32> loc(#loc98)\n %b_block_54 = arith.cmpi slt, %b_ptrs, %b_block : tensor<16x1xi32> loc(#loc98)\n %b_block_55 = tt.broadcast %b_block_54 : tensor<16x1xi1> -> tensor<16x16xi1> loc(#loc99)\n %b_block_56 = tt.load %b_ptrs_47, %b_block_55, %cst : tensor<16x16x!tt.ptr> loc(#loc99)\n %accumulator_57 = tt.dot %a_block_53, %b_block_56, %accumulator_48, inputPrecision = tf32 : tensor<16x16xf16> * tensor<16x16xf16> -> tensor<16x16xf32> loc(#loc100)\n %a_ptrs_58 = tt.addptr %a_ptrs_46, %cst_0 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc101)\n %b_ptrs_59 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc102)\n %b_ptrs_60 = tt.splat %b_ptrs_59 : i32 -> tensor<16x16xi32> loc(#loc103)\n %b_ptrs_61 = tt.addptr %b_ptrs_47, %b_ptrs_60 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc103)\n scf.yield %a_ptrs_58, %b_ptrs_61, %accumulator_57 : tensor<16x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<16x16xf32> loc(#loc43)\n } loc(#loc118)\n %c_block = arith.truncf %accumulator#2 : tensor<16x16xf32> to tensor<16x16xf16> loc(#loc104)\n %c_ptrs = tt.expand_dims %offs_am_9 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc105)\n %c_ptrs_32 = tt.splat %stride_cm : i32 -> tensor<16x1xi32> loc(#loc106)\n %c_ptrs_33 = arith.muli %c_ptrs_32, %c_ptrs : tensor<16x1xi32> loc(#loc106)\n %c_ptrs_34 = tt.splat %c : !tt.ptr -> tensor<16x1x!tt.ptr> loc(#loc107)\n %c_ptrs_35 = tt.addptr %c_ptrs_34, %c_ptrs_33 : tensor<16x1x!tt.ptr>, tensor<16x1xi32> loc(#loc107)\n %c_ptrs_36 = tt.expand_dims %offs_bn_13 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc108)\n %c_ptrs_37 = tt.broadcast %c_ptrs_35 : tensor<16x1x!tt.ptr> -> tensor<16x16x!tt.ptr> loc(#loc109)\n %c_ptrs_38 = tt.broadcast %c_ptrs_36 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc109)\n %c_ptrs_39 = tt.addptr %c_ptrs_37, %c_ptrs_38 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc109)\n %c_mask = tt.splat %M : i32 -> tensor<16x1xi32> loc(#loc110)\n %c_mask_40 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<16x1xi32> loc(#loc110)\n %c_mask_41 = tt.splat %N : i32 -> tensor<1x16xi32> loc(#loc111)\n %c_mask_42 = arith.cmpi slt, %c_ptrs_36, %c_mask_41 : tensor<1x16xi32> loc(#loc111)\n %c_mask_43 = tt.broadcast %c_mask_40 : tensor<16x1xi1> -> tensor<16x16xi1> loc(#loc112)\n %c_mask_44 = tt.broadcast %c_mask_42 : tensor<1x16xi1> -> tensor<16x16xi1> loc(#loc112)\n %c_mask_45 = arith.andi %c_mask_43, %c_mask_44 : tensor<16x16xi1> loc(#loc112)\n tt.store %c_ptrs_39, %c_block, %c_mask_45 : tensor<16x16x!tt.ptr> loc(#loc53)\n tt.return loc(#loc54)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:22)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:24)\n#loc4 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:17)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:27)\n#loc6 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:30)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:27)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:22)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:33)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:46)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:33)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:27)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:19)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:40)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:23)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:51)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:38)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:68)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:23)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:38)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:68)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:26)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:37)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:56)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:49)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:25)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:36)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:56)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:48)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:33)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:51)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:47)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:43)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:12)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:43)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:12)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:39)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:18)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:33)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:18)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:8)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:29)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:37)\n#loc46 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:29)\n#loc47 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:17)\n#loc48 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:68)\n#loc49 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:48)\n#loc50 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:33)\n#loc51 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:58)\n#loc52 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:39)\n#loc53 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:21)\n#loc54 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:4)\n#loc64 = loc(\"pid\"(#loc3))\n#loc65 = loc(\"num_pid_m\"(#loc5))\n#loc66 = loc(\"num_pid_n\"(#loc7))\n#loc67 = loc(\"group_id\"(#loc8))\n#loc68 = loc(\"group_size\"(#loc9))\n#loc69 = loc(\"group_size\"(#loc10))\n#loc70 = loc(\"pid_m\"(#loc11))\n#loc71 = loc(\"pid_m\"(#loc12))\n#loc72 = loc(\"pid_n\"(#loc13))\n#loc73 = loc(\"pid_n\"(#loc14))\n#loc74 = loc(\"offs_am\"(#loc15))\n#loc75 = loc(\"offs_am\"(#loc16))\n#loc76 = loc(\"offs_am\"(#loc17))\n#loc77 = loc(\"offs_am\"(#loc18))\n#loc78 = loc(\"offs_bn\"(#loc19))\n#loc79 = loc(\"offs_bn\"(#loc20))\n#loc80 = loc(\"offs_bn\"(#loc21))\n#loc81 = loc(\"a_ptrs\"(#loc22))\n#loc82 = loc(\"a_ptrs\"(#loc23))\n#loc83 = loc(\"a_ptrs\"(#loc24))\n#loc84 = loc(\"a_ptrs\"(#loc25))\n#loc85 = loc(\"a_ptrs\"(#loc26))\n#loc86 = loc(\"b_ptrs\"(#loc27))\n#loc87 = loc(\"b_ptrs\"(#loc28))\n#loc88 = loc(\"b_ptrs\"(#loc29))\n#loc89 = loc(\"b_ptrs\"(#loc30))\n#loc90 = loc(\"b_ptrs\"(#loc31))\n#loc91 = loc(callsite(#loc4 at #loc32))\n#loc92 = loc(callsite(#loc6 at #loc32))\n#loc93 = loc(\"a_ptrs\"(#loc2))\n#loc94 = loc(\"a_block\"(#loc33))\n#loc95 = loc(\"a_block\"(#loc34))\n#loc96 = loc(\"a_block\"(#loc35))\n#loc97 = loc(\"a_block\"(#loc36))\n#loc98 = loc(\"b_block\"(#loc37))\n#loc99 = loc(\"b_block\"(#loc38))\n#loc100 = loc(\"accumulator\"(#loc39))\n#loc101 = loc(\"a_ptrs\"(#loc40))\n#loc102 = loc(\"b_ptrs\"(#loc41))\n#loc103 = loc(\"b_ptrs\"(#loc42))\n#loc104 = loc(\"c_block\"(#loc44))\n#loc105 = loc(\"c_ptrs\"(#loc45))\n#loc106 = loc(\"c_ptrs\"(#loc46))\n#loc107 = loc(\"c_ptrs\"(#loc47))\n#loc108 = loc(\"c_ptrs\"(#loc48))\n#loc109 = loc(\"c_ptrs\"(#loc49))\n#loc110 = loc(\"c_mask\"(#loc50))\n#loc111 = loc(\"c_mask\"(#loc51))\n#loc112 = loc(\"c_mask\"(#loc52))\n#loc113 = loc(callsite(#loc4 at #loc65))\n#loc114 = loc(callsite(#loc6 at #loc65))\n#loc115 = loc(callsite(#loc4 at #loc66))\n#loc116 = loc(callsite(#loc6 at #loc66))\n#loc117 = loc(\"b_ptrs\"(#loc93))\n#loc118 = loc(\"accumulator\"(#loc117))\n","matmul_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1, 8], threadsPerWarp = [16, 2], warpsPerCTA = [1, 1], order = [1, 0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:0)\n#mma = #ttg.nvidia_mma<{versionMajor = 2, versionMinor = 0, warpsPerCTA = [1, 1], instrShape = [16, 8]}>\n#shared = #ttg.swizzled_shared<{vec = 8, perPhase = 4, maxPhase = 2, order = [1, 0]}>\n#smem = #ttg.shared_memory\n#loc55 = loc(\"a\"(#loc))\n#loc56 = loc(\"b\"(#loc))\n#loc57 = loc(\"c\"(#loc))\n#loc58 = loc(\"M\"(#loc))\n#loc59 = loc(\"N\"(#loc))\n#loc60 = loc(\"K\"(#loc))\n#loc61 = loc(\"stride_am\"(#loc))\n#loc62 = loc(\"stride_bk\"(#loc))\n#loc63 = loc(\"stride_cm\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 1 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %cst = arith.constant dense<0.000000e+00> : tensor<16x16xf32, #mma> loc(#loc1)\n %cst_0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #blocked> loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc1)\n %cst_1 = arith.constant dense<16> : tensor<16x16xi32, #blocked> loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc64)\n %num_pid_m = arith.addi %M, %c15_i32 : i32 loc(#loc113)\n %num_pid_m_2 = arith.divsi %num_pid_m, %c16_i32 : i32 loc(#loc114)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc115)\n %num_pid_n_3 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc116)\n %group_id = arith.divsi %pid, %num_pid_n_3 : i32 loc(#loc67)\n %group_size = arith.subi %num_pid_m_2, %group_id : i32 loc(#loc68)\n %group_size_4 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc69)\n %pid_m = arith.remsi %pid, %group_size_4 : i32 loc(#loc70)\n %pid_m_5 = arith.addi %group_id, %pid_m : i32 loc(#loc71)\n %pid_n = arith.remsi %pid, %num_pid_n_3 : i32 loc(#loc72)\n %pid_n_6 = arith.divsi %pid_n, %group_size_4 : i32 loc(#loc73)\n %offs_am = arith.muli %pid_m_5, %c16_i32 : i32 loc(#loc74)\n %offs_am_7 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc75)\n %offs_am_8 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc75)\n %offs_am_9 = tt.splat %offs_am : i32 -> tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc76)\n %offs_am_10 = arith.addi %offs_am_9, %offs_am_7 : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc76)\n %offs_am_11 = tt.splat %M : i32 -> tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc77)\n %offs_am_12 = arith.remsi %offs_am_10, %offs_am_11 : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc77)\n %offs_bn = arith.muli %pid_n_6, %c16_i32 : i32 loc(#loc78)\n %offs_bn_13 = tt.splat %offs_bn : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc79)\n %offs_bn_14 = arith.addi %offs_bn_13, %offs_am_8 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc79)\n %offs_bn_15 = tt.splat %N : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc80)\n %offs_bn_16 = arith.remsi %offs_bn_14, %offs_bn_15 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc80)\n %a_ptrs = tt.expand_dims %offs_am_12 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc81)\n %a_ptrs_17 = tt.splat %stride_am : i32 -> tensor<16x1xi32, #blocked> loc(#loc82)\n %a_ptrs_18 = arith.muli %a_ptrs, %a_ptrs_17 : tensor<16x1xi32, #blocked> loc(#loc82)\n %a_ptrs_19 = tt.expand_dims %offs_am_8 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc83)\n %a_ptrs_20 = tt.broadcast %a_ptrs_18 : tensor<16x1xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc84)\n %a_ptrs_21 = tt.broadcast %a_ptrs_19 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc84)\n %a_ptrs_22 = arith.addi %a_ptrs_20, %a_ptrs_21 : tensor<16x16xi32, #blocked> loc(#loc84)\n %a_ptrs_23 = tt.splat %a : !tt.ptr -> tensor<16x16x!tt.ptr, #blocked> loc(#loc85)\n %a_ptrs_24 = tt.addptr %a_ptrs_23, %a_ptrs_22 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc85)\n %b_ptrs = tt.expand_dims %offs_am_7 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc86)\n %b_ptrs_25 = tt.splat %stride_bk : i32 -> tensor<16x1xi32, #blocked> loc(#loc87)\n %b_ptrs_26 = arith.muli %b_ptrs, %b_ptrs_25 : tensor<16x1xi32, #blocked> loc(#loc87)\n %b_ptrs_27 = tt.expand_dims %offs_bn_16 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc88)\n %b_ptrs_28 = tt.broadcast %b_ptrs_26 : tensor<16x1xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc89)\n %b_ptrs_29 = tt.broadcast %b_ptrs_27 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc89)\n %b_ptrs_30 = arith.addi %b_ptrs_28, %b_ptrs_29 : tensor<16x16xi32, #blocked> loc(#loc89)\n %b_ptrs_31 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr, #blocked> loc(#loc90)\n %b_ptrs_32 = tt.addptr %b_ptrs_31, %b_ptrs_30 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc90)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc91)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc92)\n %b_ptrs_33 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc93)\n %b_ptrs_34 = tt.splat %b_ptrs_33 : i32 -> tensor<16x16xi32, #blocked> loc(#loc94)\n %accumulator:3 = scf.for %accumulator_49 = %c0_i32 to %1 step %c1_i32 iter_args(%arg10 = %cst, %a_ptrs_50 = %a_ptrs_24, %b_ptrs_51 = %b_ptrs_32) -> (tensor<16x16xf32, #mma>, tensor<16x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked>) : i32 {\n %a_block = arith.muli %accumulator_49, %c16_i32 : i32 loc(#loc96)\n %a_block_52 = arith.subi %K, %a_block : i32 loc(#loc97)\n %a_block_53 = tt.splat %a_block_52 : i32 -> tensor<1x16xi32, #blocked> loc(#loc98)\n %a_block_54 = arith.cmpi slt, %a_ptrs_19, %a_block_53 : tensor<1x16xi32, #blocked> loc(#loc98)\n %a_block_55 = tt.broadcast %a_block_54 : tensor<1x16xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc99)\n %a_block_56 = tt.load %a_ptrs_50, %a_block_55, %cst_0 : tensor<16x16x!tt.ptr, #blocked> loc(#loc99)\n %a_block_57 = ttg.local_alloc %a_block_56 : (tensor<16x16xf16, #blocked>) -> !ttg.memdesc<16x16xf16, #shared, #smem> loc(#loc99)\n %b_block = tt.splat %a_block_52 : i32 -> tensor<16x1xi32, #blocked> loc(#loc100)\n %b_block_58 = arith.cmpi slt, %b_ptrs, %b_block : tensor<16x1xi32, #blocked> loc(#loc100)\n %b_block_59 = tt.broadcast %b_block_58 : tensor<16x1xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc101)\n %b_block_60 = tt.load %b_ptrs_51, %b_block_59, %cst_0 : tensor<16x16x!tt.ptr, #blocked> loc(#loc101)\n %b_block_61 = ttg.local_alloc %b_block_60 : (tensor<16x16xf16, #blocked>) -> !ttg.memdesc<16x16xf16, #shared, #smem> loc(#loc101)\n %a_block_62 = ttg.local_load %a_block_57 : !ttg.memdesc<16x16xf16, #shared, #smem> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> loc(#loc99)\n %b_block_63 = ttg.local_load %b_block_61 : !ttg.memdesc<16x16xf16, #shared, #smem> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> loc(#loc101)\n %accumulator_64 = tt.dot %a_block_62, %b_block_63, %arg10, inputPrecision = tf32 : tensor<16x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> * tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> -> tensor<16x16xf32, #mma> loc(#loc102)\n %a_ptrs_65 = tt.addptr %a_ptrs_50, %cst_1 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc103)\n %b_ptrs_66 = tt.addptr %b_ptrs_51, %b_ptrs_34 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc94)\n scf.yield %accumulator_64, %a_ptrs_65, %b_ptrs_66 : tensor<16x16xf32, #mma>, tensor<16x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked> loc(#loc43)\n } loc(#loc118)\n %c_block = arith.truncf %accumulator#0 : tensor<16x16xf32, #mma> to tensor<16x16xf16, #mma> loc(#loc104)\n %c_ptrs = tt.expand_dims %offs_am_10 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc105)\n %c_ptrs_35 = tt.splat %stride_cm : i32 -> tensor<16x1xi32, #blocked> loc(#loc106)\n %c_ptrs_36 = arith.muli %c_ptrs_35, %c_ptrs : tensor<16x1xi32, #blocked> loc(#loc106)\n %c_ptrs_37 = tt.splat %c : !tt.ptr -> tensor<16x1x!tt.ptr, #blocked> loc(#loc107)\n %c_ptrs_38 = tt.addptr %c_ptrs_37, %c_ptrs_36 : tensor<16x1x!tt.ptr, #blocked>, tensor<16x1xi32, #blocked> loc(#loc107)\n %c_ptrs_39 = tt.expand_dims %offs_bn_14 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc108)\n %c_ptrs_40 = tt.broadcast %c_ptrs_38 : tensor<16x1x!tt.ptr, #blocked> -> tensor<16x16x!tt.ptr, #blocked> loc(#loc109)\n %c_ptrs_41 = tt.broadcast %c_ptrs_39 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc109)\n %c_ptrs_42 = tt.addptr %c_ptrs_40, %c_ptrs_41 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc109)\n %c_mask = tt.splat %M : i32 -> tensor<16x1xi32, #blocked> loc(#loc110)\n %c_mask_43 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<16x1xi32, #blocked> loc(#loc110)\n %c_mask_44 = tt.splat %N : i32 -> tensor<1x16xi32, #blocked> loc(#loc111)\n %c_mask_45 = arith.cmpi slt, %c_ptrs_39, %c_mask_44 : tensor<1x16xi32, #blocked> loc(#loc111)\n %c_mask_46 = tt.broadcast %c_mask_43 : tensor<16x1xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc112)\n %c_mask_47 = tt.broadcast %c_mask_45 : tensor<1x16xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc112)\n %c_mask_48 = arith.andi %c_mask_46, %c_mask_47 : tensor<16x16xi1, #blocked> loc(#loc112)\n %2 = ttg.convert_layout %c_block : tensor<16x16xf16, #mma> -> tensor<16x16xf16, #blocked> loc(#loc53)\n tt.store %c_ptrs_42, %2, %c_mask_48 : tensor<16x16x!tt.ptr, #blocked> loc(#loc53)\n tt.return loc(#loc54)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:24)\n#loc3 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:17)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:27)\n#loc5 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:30)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:27)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:22)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:33)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:46)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:33)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:27)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:19)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:40)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:23)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:51)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:38)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:68)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:23)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:38)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:68)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:26)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:37)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:56)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:49)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:25)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:36)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:56)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:48)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:33)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:33)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:18)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:22)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:51)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:47)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:43)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:12)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:43)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:12)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:39)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:18)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:8)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:29)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:37)\n#loc46 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:29)\n#loc47 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:17)\n#loc48 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:68)\n#loc49 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:48)\n#loc50 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:33)\n#loc51 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:58)\n#loc52 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:39)\n#loc53 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:21)\n#loc54 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:4)\n#loc64 = loc(\"pid\"(#loc2))\n#loc65 = loc(\"num_pid_m\"(#loc4))\n#loc66 = loc(\"num_pid_n\"(#loc6))\n#loc67 = loc(\"group_id\"(#loc7))\n#loc68 = loc(\"group_size\"(#loc8))\n#loc69 = loc(\"group_size\"(#loc9))\n#loc70 = loc(\"pid_m\"(#loc10))\n#loc71 = loc(\"pid_m\"(#loc11))\n#loc72 = loc(\"pid_n\"(#loc12))\n#loc73 = loc(\"pid_n\"(#loc13))\n#loc74 = loc(\"offs_am\"(#loc14))\n#loc75 = loc(\"offs_am\"(#loc15))\n#loc76 = loc(\"offs_am\"(#loc16))\n#loc77 = loc(\"offs_am\"(#loc17))\n#loc78 = loc(\"offs_bn\"(#loc18))\n#loc79 = loc(\"offs_bn\"(#loc19))\n#loc80 = loc(\"offs_bn\"(#loc20))\n#loc81 = loc(\"a_ptrs\"(#loc21))\n#loc82 = loc(\"a_ptrs\"(#loc22))\n#loc83 = loc(\"a_ptrs\"(#loc23))\n#loc84 = loc(\"a_ptrs\"(#loc24))\n#loc85 = loc(\"a_ptrs\"(#loc25))\n#loc86 = loc(\"b_ptrs\"(#loc26))\n#loc87 = loc(\"b_ptrs\"(#loc27))\n#loc88 = loc(\"b_ptrs\"(#loc28))\n#loc89 = loc(\"b_ptrs\"(#loc29))\n#loc90 = loc(\"b_ptrs\"(#loc30))\n#loc91 = loc(callsite(#loc3 at #loc31))\n#loc92 = loc(callsite(#loc5 at #loc31))\n#loc93 = loc(\"b_ptrs\"(#loc32))\n#loc94 = loc(\"b_ptrs\"(#loc33))\n#loc95 = loc(\"a_ptrs\"(#loc34))\n#loc96 = loc(\"a_block\"(#loc35))\n#loc97 = loc(\"a_block\"(#loc36))\n#loc98 = loc(\"a_block\"(#loc37))\n#loc99 = loc(\"a_block\"(#loc38))\n#loc100 = loc(\"b_block\"(#loc39))\n#loc101 = loc(\"b_block\"(#loc40))\n#loc102 = loc(\"accumulator\"(#loc41))\n#loc103 = loc(\"a_ptrs\"(#loc42))\n#loc104 = loc(\"c_block\"(#loc44))\n#loc105 = loc(\"c_ptrs\"(#loc45))\n#loc106 = loc(\"c_ptrs\"(#loc46))\n#loc107 = loc(\"c_ptrs\"(#loc47))\n#loc108 = loc(\"c_ptrs\"(#loc48))\n#loc109 = loc(\"c_ptrs\"(#loc49))\n#loc110 = loc(\"c_mask\"(#loc50))\n#loc111 = loc(\"c_mask\"(#loc51))\n#loc112 = loc(\"c_mask\"(#loc52))\n#loc113 = loc(callsite(#loc3 at #loc65))\n#loc114 = loc(callsite(#loc5 at #loc65))\n#loc115 = loc(callsite(#loc3 at #loc66))\n#loc116 = loc(callsite(#loc5 at #loc66))\n#loc117 = loc(\"b_ptrs\"(#loc95))\n#loc118 = loc(\"accumulator\"(#loc117))\n","matmul_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n@global_smem = external addrspace(3) global [0 x i8], align 16\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @matmul_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i32 %8, ptr addrspace(1) readnone captures(none) %9, ptr addrspace(1) readnone captures(none) %10) local_unnamed_addr #0 !dbg !4 {\n %12 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %13 = add i32 %3, 15, !dbg !13\n %14 = sdiv i32 %13, 16, !dbg !18\n %15 = add i32 %4, 15, !dbg !19\n %16 = sdiv i32 %15, 16, !dbg !21\n %.frozen = freeze i32 %16, !dbg !22\n %17 = sdiv i32 %12, %.frozen, !dbg !22\n %18 = sub i32 %14, %17, !dbg !23\n %19 = tail call i32 @llvm.smin.i32(i32 %18, i32 1), !dbg !24\n %20 = srem i32 %12, %19, !dbg !25\n %21 = add i32 %20, %17, !dbg !26\n %22 = mul i32 %17, %.frozen, !dbg !27\n %.decomposed = sub i32 %12, %22, !dbg !27\n %23 = sdiv i32 %.decomposed, %19, !dbg !28\n %24 = shl i32 %21, 4, !dbg !29\n %25 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !30\n %26 = lshr i32 %25, 1, !dbg !30\n %27 = and i32 %26, 15, !dbg !30\n %28 = and i32 %25, 1, !dbg !30\n %29 = shl nuw nsw i32 %28, 3, !dbg !30\n %30 = or disjoint i32 %24, %27, !dbg !31\n %31 = shl nsw i32 %23, 4, !dbg !32\n %32 = or disjoint i32 %31, %29, !dbg !33\n %33 = add i32 %5, 15, !dbg !34\n %34 = sdiv i32 %33, 16, !dbg !36\n %35 = icmp sgt i32 %33, 15, !dbg !37\n br i1 %35, label %.lr.ph, label %.._crit_edge_crit_edge, !dbg !37\n\n.._crit_edge_crit_edge: ; preds = %11\n %.pre = shl nuw nsw i32 %25, 4, !dbg !38\n %.pre65 = and i32 %25, 4, !dbg !38\n %.pre67 = and i32 %25, 8, !dbg !38\n br label %._crit_edge, !dbg !37\n\n.lr.ph: ; preds = %11\n %36 = shl i32 %7, 4, !dbg !39\n %37 = srem i32 %32, %4, !dbg !40\n %38 = mul i32 %7, %27, !dbg !41\n %39 = add i32 %37, %38, !dbg !42\n %40 = sext i32 %39 to i64, !dbg !43\n %41 = getelementptr half, ptr addrspace(1) %1, i64 %40, !dbg !43\n %42 = srem i32 %30, %3, !dbg !44\n %43 = mul i32 %42, %6, !dbg !45\n %44 = add i32 %43, %29, !dbg !46\n %45 = sext i32 %44 to i64, !dbg !47\n %46 = getelementptr half, ptr addrspace(1) %0, i64 %45, !dbg !47\n %47 = shl nuw nsw i32 %25, 4\n %48 = and i32 %47, 368\n %49 = and i32 %25, 8\n %50 = icmp eq i32 %49, 0\n %51 = select i1 %50, i32 0, i32 144\n %52 = xor i32 %51, %48\n %53 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %52\n %54 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 512), i32 %52\n %55 = shl nuw nsw i32 %25, 5\n %56 = and i32 %55, 352\n %57 = and i32 %25, 4\n %58 = icmp eq i32 %57, 0\n %59 = select i1 %58, i32 0, i32 144\n %60 = and i32 %25, 16\n %61 = xor i32 %59, %60\n %62 = or disjoint i32 %61, %56\n %63 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %62\n %64 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 512), i32 %62\n %65 = sext i32 %36 to i64\n br label %66, !dbg !37\n\n66: ; preds = %.lr.ph, %66\n %.pn3163 = phi ptr addrspace(1) [ %41, %.lr.ph ], [ %111, %66 ]\n %.pn1562 = phi ptr addrspace(1) [ %46, %.lr.ph ], [ %110, %66 ]\n %67 = phi float [ 0.000000e+00, %.lr.ph ], [ %101, %66 ]\n %68 = phi float [ 0.000000e+00, %.lr.ph ], [ %102, %66 ]\n %69 = phi float [ 0.000000e+00, %.lr.ph ], [ %103, %66 ]\n %70 = phi float [ 0.000000e+00, %.lr.ph ], [ %104, %66 ]\n %71 = phi float [ 0.000000e+00, %.lr.ph ], [ %106, %66 ]\n %72 = phi float [ 0.000000e+00, %.lr.ph ], [ %107, %66 ]\n %73 = phi float [ 0.000000e+00, %.lr.ph ], [ %108, %66 ]\n %74 = phi float [ 0.000000e+00, %.lr.ph ], [ %109, %66 ]\n %75 = phi i32 [ 0, %.lr.ph ], [ %112, %66 ]\n %76 = shl i32 %75, 4, !dbg !48\n %77 = sub i32 %5, %76, !dbg !49\n %78 = icmp slt i32 %29, %77, !dbg !50\n %79 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn1562, i1 %78) #6, !dbg !51\n %80 = extractvalue { i32, i32, i32, i32 } %79, 0, !dbg !51\n %81 = extractvalue { i32, i32, i32, i32 } %79, 1, !dbg !51\n %82 = extractvalue { i32, i32, i32, i32 } %79, 2, !dbg !51\n %83 = extractvalue { i32, i32, i32, i32 } %79, 3, !dbg !51\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !51\n tail call void asm sideeffect \"st.shared.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %53, i32 %80, i32 %81, i32 %82, i32 %83) #6, !dbg !51\n %84 = icmp slt i32 %27, %77, !dbg !52\n %85 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn3163, i1 %84) #6, !dbg !53\n %86 = extractvalue { i32, i32, i32, i32 } %85, 0, !dbg !53\n %87 = extractvalue { i32, i32, i32, i32 } %85, 1, !dbg !53\n %88 = extractvalue { i32, i32, i32, i32 } %85, 2, !dbg !53\n %89 = extractvalue { i32, i32, i32, i32 } %85, 3, !dbg !53\n tail call void asm sideeffect \"st.shared.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %54, i32 %86, i32 %87, i32 %88, i32 %89) #6, !dbg !53\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !51\n %90 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %63), !dbg !51\n %91 = extractvalue { i32, i32, i32, i32 } %90, 0, !dbg !51\n %92 = extractvalue { i32, i32, i32, i32 } %90, 1, !dbg !51\n %93 = extractvalue { i32, i32, i32, i32 } %90, 2, !dbg !51\n %94 = extractvalue { i32, i32, i32, i32 } %90, 3, !dbg !51\n %95 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %64), !dbg !53\n %96 = extractvalue { i32, i32, i32, i32 } %95, 0, !dbg !53\n %97 = extractvalue { i32, i32, i32, i32 } %95, 1, !dbg !53\n %98 = extractvalue { i32, i32, i32, i32 } %95, 2, !dbg !53\n %99 = extractvalue { i32, i32, i32, i32 } %95, 3, !dbg !53\n %100 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %67, float %68, float %69, float %70, i32 %91, i32 %92, i32 %93, i32 %94, i32 %96, i32 %97) #6, !dbg !54\n %101 = extractvalue { float, float, float, float } %100, 0, !dbg !54\n %102 = extractvalue { float, float, float, float } %100, 1, !dbg !54\n %103 = extractvalue { float, float, float, float } %100, 2, !dbg !54\n %104 = extractvalue { float, float, float, float } %100, 3, !dbg !54\n %105 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %71, float %72, float %73, float %74, i32 %91, i32 %92, i32 %93, i32 %94, i32 %98, i32 %99) #6, !dbg !54\n %106 = extractvalue { float, float, float, float } %105, 0, !dbg !54\n %107 = extractvalue { float, float, float, float } %105, 1, !dbg !54\n %108 = extractvalue { float, float, float, float } %105, 2, !dbg !54\n %109 = extractvalue { float, float, float, float } %105, 3, !dbg !54\n %110 = getelementptr i8, ptr addrspace(1) %.pn1562, i64 32, !dbg !55\n %111 = getelementptr half, ptr addrspace(1) %.pn3163, i64 %65, !dbg !56\n %112 = add nuw nsw i32 %75, 1, !dbg !37\n %exitcond.not = icmp eq i32 %112, %34, !dbg !37\n br i1 %exitcond.not, label %._crit_edge.loopexit, label %66, !dbg !37\n\n._crit_edge.loopexit: ; preds = %66\n %113 = insertelement <8 x float> poison, float %101, i64 0, !dbg !57\n %114 = insertelement <8 x float> %113, float %102, i64 1, !dbg !57\n %115 = insertelement <8 x float> %114, float %103, i64 2, !dbg !57\n %116 = insertelement <8 x float> %115, float %104, i64 3, !dbg !57\n %117 = insertelement <8 x float> %116, float %106, i64 4, !dbg !57\n %118 = insertelement <8 x float> %117, float %107, i64 5, !dbg !57\n %119 = insertelement <8 x float> %118, float %108, i64 6, !dbg !57\n %120 = insertelement <8 x float> %119, float %109, i64 7, !dbg !57\n %121 = fptrunc <8 x float> %120 to <8 x half>, !dbg !57\n %122 = bitcast <8 x half> %121 to <4 x i32>, !dbg !38\n br label %._crit_edge, !dbg !57\n\n._crit_edge: ; preds = %._crit_edge.loopexit, %.._crit_edge_crit_edge\n %.pre-phi68 = phi i32 [ %.pre67, %.._crit_edge_crit_edge ], [ %49, %._crit_edge.loopexit ], !dbg !38\n %.pre-phi66 = phi i32 [ %.pre65, %.._crit_edge_crit_edge ], [ %57, %._crit_edge.loopexit ], !dbg !38\n %.pre-phi = phi i32 [ %.pre, %.._crit_edge_crit_edge ], [ %47, %._crit_edge.loopexit ], !dbg !38\n %.bc82 = phi <4 x i32> [ zeroinitializer, %.._crit_edge_crit_edge ], [ %122, %._crit_edge.loopexit ]\n %123 = mul i32 %30, %8, !dbg !58\n %124 = sext i32 %123 to i64, !dbg !59\n %125 = getelementptr half, ptr addrspace(1) %2, i64 %124, !dbg !59\n %126 = sext i32 %32 to i64, !dbg !60\n %127 = getelementptr half, ptr addrspace(1) %125, i64 %126, !dbg !60\n %128 = icmp slt i32 %30, %3, !dbg !61\n %129 = icmp slt i32 %32, %4, !dbg !62\n %130 = and i1 %128, %129, !dbg !63\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !38\n %131 = and i32 %.pre-phi, 176, !dbg !38\n %132 = icmp eq i32 %.pre-phi66, 0, !dbg !38\n %133 = select i1 %132, i32 0, i32 320, !dbg !38\n %134 = shl nuw nsw i32 %25, 2, !dbg !38\n %135 = and i32 %134, 64, !dbg !38\n %136 = or disjoint i32 %133, %131, !dbg !38\n %137 = xor i32 %136, %135, !dbg !38\n %138 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %137, !dbg !38\n %.extract77 = extractelement <4 x i32> %.bc82, i64 0, !dbg !38\n %.extract79 = extractelement <4 x i32> %.bc82, i64 1, !dbg !38\n %.extract81 = extractelement <4 x i32> %.bc82, i64 2, !dbg !38\n %.extract83 = extractelement <4 x i32> %.bc82, i64 3, !dbg !38\n tail call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %138, i32 %.extract77, i32 %.extract79, i32 %.extract81, i32 %.extract83), !dbg !38\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !38\n %139 = shl nuw nsw i32 %25, 3, !dbg !38\n %140 = and i32 %139, 176, !dbg !38\n %141 = shl nuw nsw i32 %28, 6, !dbg !38\n %142 = icmp eq i32 %.pre-phi68, 0, !dbg !38\n %143 = select i1 %142, i32 0, i32 320, !dbg !38\n %144 = or disjoint i32 %140, %141, !dbg !38\n %145 = xor i32 %144, %143, !dbg !38\n %146 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %145, !dbg !38\n %.extract = load i32, ptr addrspace(3) %146, align 16, !dbg !38\n %147 = getelementptr inbounds nuw i8, ptr addrspace(3) %146, i32 4, !dbg !38\n %.extract57 = load i32, ptr addrspace(3) %147, align 4, !dbg !38\n %148 = getelementptr inbounds nuw i8, ptr addrspace(3) %146, i32 8, !dbg !38\n %.extract59 = load i32, ptr addrspace(3) %148, align 8, !dbg !38\n %149 = getelementptr inbounds nuw i8, ptr addrspace(3) %146, i32 12, !dbg !38\n %.extract61 = load i32, ptr addrspace(3) %149, align 4, !dbg !38\n tail call void asm sideeffect \"@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };\", \"r,r,r,r,l,b\"(i32 %.extract, i32 %.extract57, i32 %.extract59, i32 %.extract61, ptr addrspace(1) %127, i1 %130) #6, !dbg !38\n ret void, !dbg !64\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)\ndeclare i32 @llvm.smin.i32(i32, i32) #2\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3\n\n; Function Attrs: convergent nocallback nounwind memory(argmem: write)\ndeclare void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) writeonly captures(none), i32, i32, i32, i32) #4\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.bar.warp.sync(i32) #3\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"32\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #3 = { convergent nocallback nounwind }\nattributes #4 = { convergent nocallback nounwind memory(argmem: write) }\nattributes #5 = { convergent nocallback nofree nounwind memory(argmem: read) }\nattributes #6 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"matmul_kernel\", linkageName: \"matmul_kernel\", scope: !1, file: !1, line: 292, type: !5, scopeLine: 292, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !9, !9, !9, !9, !9, !9, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"half\", size: 16, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 311, column: 24, scope: !4)\n!13 = !DILocation(line: 43, column: 17, scope: !14, inlinedAt: !16)\n!14 = distinct !DILexicalBlockFile(scope: !4, file: !15, discriminator: 0)\n!15 = !DIFile(filename: \"standard.py\", directory: \"/python/lib/python3.14/site-packages/triton/language\")\n!16 = !DILocation(line: 312, column: 27, scope: !17)\n!17 = distinct !DILexicalBlockFile(scope: !4, file: !1, discriminator: 0)\n!18 = !DILocation(line: 43, column: 30, scope: !14, inlinedAt: !16)\n!19 = !DILocation(line: 43, column: 17, scope: !14, inlinedAt: !20)\n!20 = !DILocation(line: 313, column: 27, scope: !17)\n!21 = !DILocation(line: 43, column: 30, scope: !14, inlinedAt: !20)\n!22 = !DILocation(line: 315, column: 22, scope: !4)\n!23 = !DILocation(line: 317, column: 33, scope: !4)\n!24 = !DILocation(line: 317, column: 46, scope: !4)\n!25 = !DILocation(line: 318, column: 33, scope: !4)\n!26 = !DILocation(line: 318, column: 27, scope: !4)\n!27 = !DILocation(line: 319, column: 19, scope: !4)\n!28 = !DILocation(line: 319, column: 40, scope: !4)\n!29 = !DILocation(line: 321, column: 23, scope: !4)\n!30 = !DILocation(line: 321, column: 51, scope: !4)\n!31 = !DILocation(line: 321, column: 38, scope: !4)\n!32 = !DILocation(line: 322, column: 23, scope: !4)\n!33 = !DILocation(line: 322, column: 38, scope: !4)\n!34 = !DILocation(line: 43, column: 17, scope: !14, inlinedAt: !35)\n!35 = !DILocation(line: 328, column: 33, scope: !17)\n!36 = !DILocation(line: 43, column: 30, scope: !14, inlinedAt: !35)\n!37 = !DILocation(line: 328, column: 22, scope: !4)\n!38 = !DILocation(line: 344, column: 21, scope: !4)\n!39 = !DILocation(line: 337, column: 33, scope: !4)\n!40 = !DILocation(line: 322, column: 68, scope: !4)\n!41 = !DILocation(line: 325, column: 36, scope: !4)\n!42 = !DILocation(line: 325, column: 48, scope: !4)\n!43 = !DILocation(line: 325, column: 18, scope: !4)\n!44 = !DILocation(line: 321, column: 68, scope: !4)\n!45 = !DILocation(line: 324, column: 37, scope: !4)\n!46 = !DILocation(line: 324, column: 49, scope: !4)\n!47 = !DILocation(line: 324, column: 18, scope: !4)\n!48 = !DILocation(line: 330, column: 51, scope: !4)\n!49 = !DILocation(line: 330, column: 47, scope: !4)\n!50 = !DILocation(line: 330, column: 43, scope: !4)\n!51 = !DILocation(line: 330, column: 12, scope: !4)\n!52 = !DILocation(line: 333, column: 43, scope: !4)\n!53 = !DILocation(line: 333, column: 12, scope: !4)\n!54 = !DILocation(line: 335, column: 39, scope: !4)\n!55 = !DILocation(line: 336, column: 18, scope: !4)\n!56 = !DILocation(line: 337, column: 18, scope: !4)\n!57 = !DILocation(line: 338, column: 29, scope: !4)\n!58 = !DILocation(line: 342, column: 29, scope: !4)\n!59 = !DILocation(line: 342, column: 17, scope: !4)\n!60 = !DILocation(line: 342, column: 48, scope: !4)\n!61 = !DILocation(line: 343, column: 33, scope: !4)\n!62 = !DILocation(line: 343, column: 58, scope: !4)\n!63 = !DILocation(line: 343, column: 39, scope: !4)\n!64 = !DILocation(line: 344, column: 4, scope: !4)\n","matmul_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.7\n.target sm_90a\n.address_size 64\n\n\t// .globl\tmatmul_kernel // -- Begin function matmul_kernel\n.extern .shared .align 16 .b8 global_smem[];\n // @matmul_kernel\n.visible .entry matmul_kernel(\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_2,\n\t.param .u32 matmul_kernel_param_3,\n\t.param .u32 matmul_kernel_param_4,\n\t.param .u32 matmul_kernel_param_5,\n\t.param .u32 matmul_kernel_param_6,\n\t.param .u32 matmul_kernel_param_7,\n\t.param .u32 matmul_kernel_param_8,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_9,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_10\n)\n.reqntid 32\n{\n\t.reg .pred \t%p<10>;\n\t.reg .b32 \t%r<111>;\n\t.reg .b64 \t%rd<9>;\n\t.loc\t1 292 0 // generate_examples.py:292:0\n$L__func_begin0:\n\t.loc\t1 292 0 // generate_examples.py:292:0\n\n// %bb.0:\n\tld.param.b32 \t%r13, [matmul_kernel_param_8];\n\tld.param.b32 \t%r94, [matmul_kernel_param_5];\n\tld.param.b32 \t%r10, [matmul_kernel_param_4];\n\tld.param.b32 \t%r9, [matmul_kernel_param_3];\n\tld.param.b64 \t%rd4, [matmul_kernel_param_2];\n$L__tmp0:\n\t.loc\t1 311 24 // generate_examples.py:311:24\n\tmov.u32 \t%r14, %ctaid.x;\n$L__tmp1:\n\t.loc\t2 43 17 // standard.py:43:17 @[ generate_examples.py:312:27 ]\n\tadd.s32 \t%r15, %r9, 15;\n\t.loc\t2 43 30 // standard.py:43:30 @[ generate_examples.py:312:27 ]\n\tshr.s32 \t%r16, %r15, 31;\n\tshr.u32 \t%r17, %r16, 28;\n\tadd.s32 \t%r18, %r15, %r17;\n\tshr.s32 \t%r19, %r18, 4;\n$L__tmp2:\n\t.loc\t2 43 17 // standard.py:43:17 @[ generate_examples.py:313:27 ]\n\tadd.s32 \t%r20, %r10, 15;\n\t.loc\t2 43 30 // standard.py:43:30 @[ generate_examples.py:313:27 ]\n\tshr.s32 \t%r21, %r20, 31;\n\tshr.u32 \t%r22, %r21, 28;\n\tadd.s32 \t%r23, %r20, %r22;\n\tshr.s32 \t%r24, %r23, 4;\n$L__tmp3:\n\t.loc\t1 315 22 // generate_examples.py:315:22\n\tdiv.s32 \t%r25, %r14, %r24;\n\t.loc\t1 317 33 // generate_examples.py:317:33\n\tsub.s32 \t%r26, %r19, %r25;\n\t.loc\t1 317 46 // generate_examples.py:317:46\n\tmin.s32 \t%r27, %r26, 1;\n\t.loc\t1 318 33 // generate_examples.py:318:33\n\trem.s32 \t%r28, %r14, %r27;\n\t.loc\t1 318 27 // generate_examples.py:318:27\n\tadd.s32 \t%r29, %r28, %r25;\n\t.loc\t1 319 19 // generate_examples.py:319:19\n\tmul.lo.s32 \t%r30, %r25, %r24;\n\tsub.s32 \t%r31, %r14, %r30;\n\t.loc\t1 319 40 // generate_examples.py:319:40\n\tdiv.s32 \t%r32, %r31, %r27;\n\t.loc\t1 321 23 // generate_examples.py:321:23\n\tshl.b32 \t%r33, %r29, 4;\n\t.loc\t1 321 51 // generate_examples.py:321:51\n\tmov.u32 \t%r1, %tid.x;\n\tbfe.u32 \t%r2, %r1, 1, 4;\n\tand.b32 \t%r3, %r1, 1;\n\tshl.b32 \t%r4, %r3, 3;\n\t.loc\t1 321 38 // generate_examples.py:321:38\n\tor.b32 \t%r5, %r33, %r2;\n\t.loc\t1 322 23 // generate_examples.py:322:23\n\tshl.b32 \t%r34, %r32, 4;\n\t.loc\t1 322 38 // generate_examples.py:322:38\n\tor.b32 \t%r6, %r34, %r4;\n$L__tmp4:\n\t.loc\t2 43 17 // standard.py:43:17 @[ generate_examples.py:328:33 ]\n\tadd.s32 \t%r35, %r94, 15;\n$L__tmp5:\n\t.loc\t1 328 22 // generate_examples.py:328:22\n\tsetp.gt.s32 \t%p1, %r35, 15;\n\tmov.b32 \t%r93, global_smem;\n\t@%p1 bra \t$L__BB0_2;\n\tbra.uni \t$L__BB0_1;\n$L__BB0_2: // %.lr.ph\n\t.loc\t1 0 22 // generate_examples.py:0:22\n\tld.param.b32 \t%r12, [matmul_kernel_param_7];\n\tld.param.b32 \t%r11, [matmul_kernel_param_6];\n\tld.param.b64 \t%rd3, [matmul_kernel_param_1];\n\tld.param.b64 \t%rd2, [matmul_kernel_param_0];\n\tshr.s32 \t%r36, %r35, 31;\n\tshr.u32 \t%r37, %r36, 28;\n\tadd.s32 \t%r38, %r35, %r37;\n\tshr.s32 \t%r95, %r38, 4;\n\t.loc\t1 337 33 // generate_examples.py:337:33\n\tshl.b32 \t%r39, %r12, 4;\n\t.loc\t1 322 68 // generate_examples.py:322:68\n\trem.s32 \t%r40, %r6, %r10;\n\t.loc\t1 325 48 // generate_examples.py:325:48\n\tmad.lo.s32 \t%r41, %r12, %r2, %r40;\n\t.loc\t1 325 18 // generate_examples.py:325:18\n\tmad.wide.s32 \t%rd7, %r41, 2, %rd3;\n\t.loc\t1 321 68 // generate_examples.py:321:68\n\trem.s32 \t%r42, %r5, %r9;\n\t.loc\t1 324 49 // generate_examples.py:324:49\n\tmad.lo.s32 \t%r43, %r42, %r11, %r4;\n\t.loc\t1 324 18 // generate_examples.py:324:18\n\tmad.wide.s32 \t%rd8, %r43, 2, %rd2;\n\tshl.b32 \t%r106, %r1, 4;\n\tand.b32 \t%r44, %r106, 368;\n\tand.b32 \t%r104, %r1, 8;\n\tbfe.s32 \t%r45, %r1, 3, 1;\n\tand.b32 \t%r46, %r45, 144;\n\txor.b32 \t%r47, %r46, %r44;\n\tadd.s32 \t%r60, %r93, %r47;\n\tadd.s32 \t%r65, %r60, 512;\n\tshl.b32 \t%r48, %r1, 5;\n\tand.b32 \t%r49, %r48, 352;\n\tand.b32 \t%r105, %r1, 4;\n\tbfe.s32 \t%r50, %r1, 2, 1;\n\tand.b32 \t%r51, %r50, 144;\n\tand.b32 \t%r52, %r1, 16;\n\txor.b32 \t%r53, %r51, %r52;\n\tor.b32 \t%r54, %r53, %r49;\n\tadd.s32 \t%r7, %r93, %r54;\n\tadd.s32 \t%r8, %r7, 512;\n\t.loc\t1 328 22 // generate_examples.py:328:22\n\tmul.wide.s32 \t%rd1, %r39, 2;\n\tmov.b32 \t%r96, 0f00000000;\n\tmov.b32 \t%r97, %r96;\n\tmov.b32 \t%r98, %r96;\n\tmov.b32 \t%r99, %r96;\n\tmov.b32 \t%r100, %r96;\n\tmov.b32 \t%r101, %r96;\n\tmov.b32 \t%r102, %r96;\n\tmov.b32 \t%r103, %r96;\n$L__BB0_3: // =>This Inner Loop Header: Depth=1\n\t.loc\t1 330 43 // generate_examples.py:330:43\n\tsetp.lt.s32 \t%p2, %r4, %r94;\n\tmov.b32 \t%r59, 0;\n\t.loc\t1 330 12 // generate_examples.py:330:12\n\t// begin inline asm\n\tmov.u32 %r55, %r59;\n\tmov.u32 %r56, %r59;\n\tmov.u32 %r57, %r59;\n\tmov.u32 %r58, %r59;\n\t@%p2 ld.global.v4.b32 { %r55, %r56, %r57, %r58 }, [ %rd8 + 0 ];\n\t// end inline asm\n\tbar.sync \t0;\n\t// begin inline asm\n\tst.shared.v4.b32 [ %r60 + 0 ], { %r55, %r56, %r57, %r58 };\n\t// end inline asm\n\t.loc\t1 333 43 // generate_examples.py:333:43\n\tsetp.lt.s32 \t%p3, %r2, %r94;\n\t.loc\t1 333 12 // generate_examples.py:333:12\n\t// begin inline asm\n\tmov.u32 %r61, %r59;\n\tmov.u32 %r62, %r59;\n\tmov.u32 %r63, %r59;\n\tmov.u32 %r64, %r59;\n\t@%p3 ld.global.v4.b32 { %r61, %r62, %r63, %r64 }, [ %rd7 + 0 ];\n\t// end inline asm\n\t// begin inline asm\n\tst.shared.v4.b32 [ %r65 + 0 ], { %r61, %r62, %r63, %r64 };\n\t// end inline asm\n\t.loc\t1 330 12 // generate_examples.py:330:12\n\tbar.sync \t0;\n\tldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r66, %r67, %r68, %r69}, [%r7];\n\t.loc\t1 333 12 // generate_examples.py:333:12\n\tldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r70, %r71, %r72, %r73}, [%r8];\n\t.loc\t1 335 39 // generate_examples.py:335:39\n\t// begin inline asm\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r96, %r97, %r98, %r99 }, { %r66, %r67, %r68, %r69 }, { %r70, %r71 }, { %r96, %r97, %r98, %r99 };\n\t// end inline asm\n\t// begin inline asm\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r100, %r101, %r102, %r103 }, { %r66, %r67, %r68, %r69 }, { %r72, %r73 }, { %r100, %r101, %r102, %r103 };\n\t// end inline asm\n\t.loc\t1 336 18 // generate_examples.py:336:18\n\tadd.s64 \t%rd8, %rd8, 32;\n\t.loc\t1 328 22 // generate_examples.py:328:22\n\tadd.s64 \t%rd7, %rd7, %rd1;\n\tadd.s32 \t%r95, %r95, -1;\n\tadd.s32 \t%r94, %r94, -16;\n\tsetp.ne.b32 \t%p4, %r95, 0;\n\t@%p4 bra \t$L__BB0_3;\n// %bb.4: // %._crit_edge.loopexit\n\t.loc\t1 338 29 // generate_examples.py:338:29\n\tcvt.rn.f16x2.f32 \t%r110, %r103, %r102;\n\tcvt.rn.f16x2.f32 \t%r109, %r101, %r100;\n\tcvt.rn.f16x2.f32 \t%r108, %r99, %r98;\n\tcvt.rn.f16x2.f32 \t%r107, %r97, %r96;\n\tbra.uni \t$L__BB0_5;\n$L__BB0_1: // %.._crit_edge_crit_edge\n\t.loc\t1 344 21 // generate_examples.py:344:21\n\tshl.b32 \t%r106, %r1, 4;\n\tand.b32 \t%r105, %r1, 4;\n\tand.b32 \t%r104, %r1, 8;\n\tmov.b32 \t%r107, 0;\n\tmov.b32 \t%r108, %r107;\n\tmov.b32 \t%r109, %r107;\n\tmov.b32 \t%r110, %r107;\n$L__BB0_5: // %._crit_edge\n\t.loc\t1 342 29 // generate_examples.py:342:29\n\tmul.lo.s32 \t%r78, %r5, %r13;\n\t.loc\t1 342 17 // generate_examples.py:342:17\n\tmad.wide.s32 \t%rd6, %r78, 2, %rd4;\n\t.loc\t1 342 48 // generate_examples.py:342:48\n\tmad.wide.s32 \t%rd5, %r6, 2, %rd6;\n\t.loc\t1 343 33 // generate_examples.py:343:33\n\tsetp.lt.s32 \t%p6, %r5, %r9;\n\t.loc\t1 343 58 // generate_examples.py:343:58\n\tsetp.lt.s32 \t%p7, %r6, %r10;\n\t.loc\t1 343 39 // generate_examples.py:343:39\n\tand.pred \t%p5, %p6, %p7;\n\t.loc\t1 344 21 // generate_examples.py:344:21\n\tbar.sync \t0;\n\tand.b32 \t%r79, %r106, 176;\n\tsetp.eq.b32 \t%p8, %r105, 0;\n\tselp.b32 \t%r80, 0, 320, %p8;\n\tshl.b32 \t%r81, %r1, 2;\n\tand.b32 \t%r82, %r81, 64;\n\tor.b32 \t%r83, %r80, %r79;\n\txor.b32 \t%r84, %r83, %r82;\n\tadd.s32 \t%r85, %r93, %r84;\n\tstmatrix.sync.aligned.m8n8.x4.shared.b16 [%r85], {%r107, %r108, %r109, %r110};\n\tbar.warp.sync \t-1;\n\tshl.b32 \t%r86, %r1, 3;\n\tand.b32 \t%r87, %r86, 176;\n\tshl.b32 \t%r88, %r3, 6;\n\tsetp.eq.b32 \t%p9, %r104, 0;\n\tselp.b32 \t%r89, 0, 320, %p9;\n\tor.b32 \t%r90, %r87, %r88;\n\txor.b32 \t%r91, %r90, %r89;\n\tadd.s32 \t%r92, %r93, %r91;\n\tld.shared.v4.b32 \t{%r74, %r75, %r76, %r77}, [%r92];\n\t// begin inline asm\n\t@%p5 st.global.v4.b32 [ %rd5 + 0 ], { %r74, %r75, %r76, %r77 };\n\t// end inline asm\n\t.loc\t1 344 4 // generate_examples.py:344:4\n\tret;\n$L__tmp6:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.file\t2 \"/python/lib/python3.14/site-packages/triton/language/standard.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 1 // DW_CHILDREN_yes\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 2 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 0 // DW_CHILDREN_no\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 32 // DW_AT_inline\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 3 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 1 // DW_CHILDREN_yes\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 4 // Abbreviation Code\n.b8 29 // DW_TAG_inlined_subroutine\n.b8 0 // DW_CHILDREN_no\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 88 // DW_AT_call_file\n.b8 11 // DW_FORM_data1\n.b8 89 // DW_AT_call_line\n.b8 5 // DW_FORM_data2\n.b8 87 // DW_AT_call_column\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 245 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0xee DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 45\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 45\n.b8 114\n.b8 101\n.b8 103\n.b8 101\n.b8 110\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n.b8 2 // Abbrev [2] 0x87:0x10 DW_TAG_subprogram\n.b8 109 // DW_AT_name\n.b8 97\n.b8 116\n.b8 109\n.b8 117\n.b8 108\n.b8 95\n.b8 107\n.b8 101\n.b8 114\n.b8 110\n.b8 101\n.b8 108\n.b8 0\n.b8 1 // DW_AT_inline\n.b8 3 // Abbrev [3] 0x97:0x61 DW_TAG_subprogram\n.b64 $L__func_begin0 // DW_AT_low_pc\n.b64 $L__func_end0 // DW_AT_high_pc\n.b32 135 // DW_AT_abstract_origin\n.b8 4 // Abbrev [4] 0xac:0x19 DW_TAG_inlined_subroutine\n.b32 135 // DW_AT_abstract_origin\n.b64 $L__tmp1 // DW_AT_low_pc\n.b64 $L__tmp2 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 56 // DW_AT_call_line\n.b8 1\n.b8 27 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xc5:0x19 DW_TAG_inlined_subroutine\n.b32 135 // DW_AT_abstract_origin\n.b64 $L__tmp2 // DW_AT_low_pc\n.b64 $L__tmp3 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 57 // DW_AT_call_line\n.b8 1\n.b8 27 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xde:0x19 DW_TAG_inlined_subroutine\n.b32 135 // DW_AT_abstract_origin\n.b64 $L__tmp4 // DW_AT_low_pc\n.b64 $L__tmp5 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 72 // DW_AT_call_line\n.b8 1\n.b8 33 // DW_AT_call_column\n.b8 0 // End Of Children Mark\n.b8 0 // End Of Children Mark\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","matmul_kernel.json":"{\"hash\": \"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 1, \"num_ctas\": 1, \"num_stages\": 1, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.7.0\", \"tensordesc_meta\": [], \"shared\": 1024, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"matmul_kernel\"}","matmul_kernel.sass":"Function:matmul_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*0010*/ LDC.64 R16, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 311\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*0020*/ S2R R7, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 94\n /*0030*/ UMOV UR4, 0x400 ;\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0040*/ ULDC.64 UR6, c[0x0][0x208] ;\n\t//## File \".nv_debug_ptx_txt\", line 94\n /*0050*/ S2UR UR5, SR_CgaCtaId ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 53\n /*0060*/ VIADD R0, R17, 0xf ;\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*0070*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 94\n /*0080*/ ULEA UR4, UR5, UR4, 0x18 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 57\n /*0090*/ LEA.HI R3, R3, R0, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*00a0*/ IABS R8, R7 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 58\n /*00b0*/ SHF.R.S32.HI R4, RZ, 0x4, R3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*00c0*/ IABS R5, R4.reuse ;\n /*00d0*/ IABS R10, R4 ;\n /*00e0*/ I2F.RP R0, R5 ;\n /*00f0*/ MUFU.RCP R0, R0 ;\n /*0100*/ VIADD R2, R0, 0xffffffe ;\n /*0110*/ IMAD.MOV R0, RZ, RZ, -R10 ;\n /*0120*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0130*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0140*/ IMAD.MOV R6, RZ, RZ, -R3 ;\n /*0150*/ IMAD R9, R6, R5, RZ ;\n /*0160*/ IMAD.MOV.U32 R6, RZ, RZ, R8 ;\n /*0170*/ IMAD.HI.U32 R3, R3, R9, R2 ;\n /*0180*/ IMAD.HI.U32 R3, R3, R6, RZ ;\n /*0190*/ IMAD R0, R3, R0, R6 ;\n /*01a0*/ ISETP.GT.U32.AND P1, PT, R5, R0, PT ;\n /*01b0*/ @!P1 IMAD.IADD R0, R0, 0x1, -R5 ;\n /*01c0*/ @!P1 VIADD R3, R3, 0x1 ;\n /*01d0*/ ISETP.NE.AND P1, PT, R4, RZ, PT ;\n /*01e0*/ ISETP.GE.U32.AND P0, PT, R0, R5, PT ;\n /*01f0*/ LOP3.LUT R0, R7, R4, RZ, 0x3c, !PT ;\n /*0200*/ ISETP.GE.AND P2, PT, R0, RZ, PT ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*0210*/ VIADD R0, R16, 0xf ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*0220*/ @P0 VIADD R3, R3, 0x1 ;\n /*0230*/ IMAD.MOV.U32 R10, RZ, RZ, R3 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 47\n /*0240*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*0250*/ @!P2 IMAD.MOV R10, RZ, RZ, -R10 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 49\n /*0260*/ LEA.HI R3, R3, R0, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*0270*/ @!P1 LOP3.LUT R10, RZ, R4, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0280*/ LEA.HI.SX32 R3, R3, -R10, 0x1c ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 72\n /*0290*/ IMAD.MOV R5, RZ, RZ, -R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*02a0*/ VIMNMX R11, R3, 0x1, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 72\n /*02b0*/ IMAD R4, R4, R5, R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*02c0*/ IABS R9, R11.reuse ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*02d0*/ IABS R12, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*02e0*/ I2F.RP R0, R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*02f0*/ LOP3.LUT R4, R4, R11, RZ, 0x3c, !PT ;\n /*0300*/ ISETP.GE.AND P2, PT, R4, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0310*/ MUFU.RCP R0, R0 ;\n /*0320*/ VIADD R2, R0, 0xffffffe ;\n /*0330*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0340*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0350*/ IMAD.MOV R8, RZ, RZ, -R3 ;\n /*0360*/ IMAD R5, R8, R9, RZ ;\n /*0370*/ IABS R8, R11 ;\n /*0380*/ IMAD.HI.U32 R3, R3, R5, R2 ;\n /*0390*/ IMAD.MOV R5, RZ, RZ, -R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*03a0*/ IMAD.MOV.U32 R8, RZ, RZ, R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*03b0*/ IMAD.HI.U32 R0, R3, R6, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*03c0*/ IMAD.HI.U32 R3, R3, R8, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*03d0*/ IMAD R2, R0, R5.reuse, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*03e0*/ IMAD R5, R3, R5, R8 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 90\n /*03f0*/ LDC R6, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*0400*/ S2R R0, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0410*/ ISETP.GT.U32.AND P1, PT, R9.reuse, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0420*/ ISETP.GT.U32.AND P0, PT, R9, R5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0430*/ @!P1 IMAD.IADD R2, R2, 0x1, -R9.reuse ;\n /*0440*/ ISETP.GE.AND P1, PT, R7, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0450*/ @!P0 IMAD.IADD R5, R5, 0x1, -R9 ;\n /*0460*/ @!P0 VIADD R3, R3, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0470*/ ISETP.GT.U32.AND P4, PT, R9, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0480*/ ISETP.GE.U32.AND P3, PT, R5, R9, PT ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 90\n /*0490*/ VIADD R4, R6, 0xf ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*04a0*/ ISETP.NE.AND P0, PT, R11, RZ, PT ;\n /*04b0*/ @!P4 IMAD.IADD R2, R2, 0x1, -R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*04c0*/ @P3 VIADD R3, R3, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 93\n /*04d0*/ ISETP.GT.AND P3, PT, R4, 0xf, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*04e0*/ @!P1 IMAD.MOV R2, RZ, RZ, -R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*04f0*/ IMAD.MOV.U32 R6, RZ, RZ, R3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0500*/ LOP3.LUT R3, RZ, R11, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0510*/ @!P2 IMAD.MOV R6, RZ, RZ, -R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0520*/ SEL R5, R3.reuse, R2, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 80\n /*0530*/ LOP3.LUT R2, R0, 0x1, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0540*/ SEL R7, R3, R6, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0550*/ IMAD.IADD R5, R10, 0x1, R5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0560*/ SHF.R.U32.HI R3, RZ, 0x1, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*0570*/ IMAD.SHL.U32 R12, R2, 0x8, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 208\n /*0580*/ @!P3 LOP3.LUT R41, R0.reuse, 0x4, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0590*/ IMAD.SHL.U32 R6, R5, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*05a0*/ SGXT.U32 R3, R3, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 207\n /*05b0*/ @!P3 IMAD.SHL.U32 R15, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 209\n /*05c0*/ @!P3 LOP3.LUT R40, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 213\n /*05d0*/ @!P3 CS2R R8, SRZ ;\n /*05e0*/ @!P3 CS2R R10, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 87\n /*05f0*/ IMAD R14, R7, 0x10, R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 83\n /*0600*/ LOP3.LUT R13, R6, R3, RZ, 0xfc, !PT ;\n /*0610*/ @!P3 BRA `(.L_x_0) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0620*/ IABS R18, R17 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0630*/ LDC R37, c[0x0][0x238] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0640*/ IABS R20, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0650*/ ULDC UR5, c[0x0][0x234] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0660*/ I2F.RP R5, R18 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0670*/ ISETP.GE.AND P2, PT, R13, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0680*/ ULDC.64 UR8, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0690*/ ISETP.NE.AND P4, PT, R16, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*06a0*/ ULDC.64 UR10, c[0x0][0x210] ;\n\t//## File \".nv_debug_ptx_txt\", line 123\n /*06b0*/ LOP3.LUT R40, R0.reuse, 0x8, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 131\n /*06c0*/ LOP3.LUT R41, R0, 0x4, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*06d0*/ I2F.RP R10, R20 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*06e0*/ MUFU.RCP R5, R5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*06f0*/ MUFU.RCP R10, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0700*/ VIADD R6, R5, 0xffffffe ;\n /*0710*/ IABS R5, R14 ;\n /*0720*/ F2I.FTZ.U32.TRUNC.NTZ R7, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0730*/ VIADD R8, R10, 0xffffffe ;\n /*0740*/ F2I.FTZ.U32.TRUNC.NTZ R9, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0750*/ IMAD.MOV.U32 R6, RZ, RZ, RZ ;\n /*0760*/ IMAD.MOV R11, RZ, RZ, -R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0770*/ IMAD.MOV.U32 R8, RZ, RZ, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0780*/ IMAD R11, R11, R18, RZ ;\n /*0790*/ IMAD.HI.U32 R6, R7, R11, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*07a0*/ IABS R7, R13 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*07b0*/ SHF.R.S32.HI R11, RZ, 0x3, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*07c0*/ IMAD.MOV R15, RZ, RZ, -R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*07d0*/ IMAD.HI.U32 R6, R6, R5, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*07e0*/ SGXT R11, R11, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*07f0*/ IMAD R15, R15, R20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0800*/ IMAD.MOV R6, RZ, RZ, -R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0810*/ IMAD.HI.U32 R8, R9, R15, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 132\n /*0820*/ SHF.R.S32.HI R9, RZ, 0x2, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0830*/ IMAD R5, R18, R6, R5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 132\n /*0840*/ SGXT R9, R9, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0850*/ IMAD.HI.U32 R8, R8, R7, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0860*/ ISETP.GT.U32.AND P0, PT, R18, R5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0870*/ IMAD.MOV R8, RZ, RZ, -R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 133\n /*0880*/ LOP3.LUT R9, R9, 0x90, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 121\n /*0890*/ IMAD.SHL.U32 R15, R0, 0x10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*08a0*/ IMAD R7, R20, R8, R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 135\n /*08b0*/ LOP3.LUT R9, R9, 0x10, R0, 0x78, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 129\n /*08c0*/ IMAD.SHL.U32 R8, R0, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*08d0*/ ISETP.GT.U32.AND P1, PT, R20, R7, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 136\n /*08e0*/ LOP3.LUT R36, R9, 0x160, R8, 0xf8, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*08f0*/ @!P0 IMAD.IADD R5, R5, 0x1, -R18 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 148\n /*0900*/ CS2R R8, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0910*/ ISETP.GT.U32.AND P3, PT, R18, R5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0920*/ @!P1 IMAD.IADD R7, R7, 0x1, -R20 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0930*/ ISETP.GE.AND P1, PT, R14, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0940*/ ISETP.GT.U32.AND P0, PT, R20, R7, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0950*/ @!P3 IMAD.IADD R5, R5, 0x1, -R18 ;\n /*0960*/ ISETP.NE.AND P3, PT, R17, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 125\n /*0970*/ LOP3.LUT R18, R11, 0x90, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0980*/ IMAD.MOV.U32 R6, RZ, RZ, R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 103\n /*0990*/ SHF.R.S32.HI R5, RZ, 0x1f, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 148\n /*09a0*/ CS2R R10, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*09b0*/ LOP3.LUT R35, R18, 0x170, R15, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*09c0*/ @!P1 IMAD.MOV R6, RZ, RZ, -R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*09d0*/ @!P0 IMAD.IADD R7, R7, 0x1, -R20 ;\n /*09e0*/ @!P2 IMAD.MOV R7, RZ, RZ, -R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*09f0*/ @!P3 LOP3.LUT R6, RZ, R17, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0a00*/ @!P4 LOP3.LUT R7, RZ, R16, RZ, 0x33, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 105\n /*0a10*/ LEA.HI R16, R5, R4, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0a20*/ IMAD R6, R3, R37, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 148\n /*0a30*/ CS2R R4, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0a40*/ IMAD R7, R7, UR5, R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 106\n /*0a50*/ SHF.R.S32.HI R34, RZ, 0x4, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 337\n\t//## File \".nv_debug_ptx_txt\", line 108\n /*0a60*/ IMAD.SHL.U32 R37, R37, 0x10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0a70*/ LEA R24, P0, R6, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 39\n /*0a80*/ ULDC UR5, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0a90*/ LEA R38, P1, R7, UR10, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0aa0*/ LEA.HI.X.SX32 R25, R6, UR9, 0x1, P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0ab0*/ LEA.HI.X.SX32 R39, R7, UR11, 0x1, P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 148\n /*0ac0*/ CS2R R6, SRZ ;\n.L_x_1:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 151\n /*0ad0*/ ISETP.GE.AND P0, PT, R12, UR5, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 159\n /*0ae0*/ CS2R R16, SRZ ;\n /*0af0*/ CS2R R18, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 166\n /*0b00*/ ISETP.GE.AND P1, PT, R3, UR5, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 173\n /*0b10*/ CS2R R20, SRZ ;\n /*0b20*/ CS2R R22, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 159\n /*0b30*/ @!P0 IMAD.MOV.U32 R32, RZ, RZ, R38 ;\n /*0b40*/ @!P0 IMAD.MOV.U32 R33, RZ, RZ, R39 ;\n /*0b50*/ @!P0 LDG.E.128 R16, desc[UR6][R32.64] ;\n\t//## File \".nv_debug_ptx_txt\", line 161\n /*0b60*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n /*0b70*/ IMAD.MOV.U32 R32, RZ, RZ, R24 ;\n /*0b80*/ IMAD.MOV.U32 R33, RZ, RZ, R25 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 173\n /*0b90*/ @!P1 LDG.E.128 R20, desc[UR6][R32.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 194\n /*0ba0*/ VIADD R34, R34, 0xffffffff ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 191\n /*0bb0*/ IADD3 R38, P1, R38, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 195\n /*0bc0*/ UIADD3 UR5, UR5, -0x10, URZ ;\n\t//## File \".nv_debug_ptx_txt\", line 196\n /*0bd0*/ ISETP.NE.U32.AND P0, PT, R34, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 191\n /*0be0*/ IMAD.X R39, RZ, RZ, R39, P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 163\n /*0bf0*/ STS.128 [R35+UR4], R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 176\n /*0c00*/ STS.128 [R35+UR4+0x200], R20 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 179\n /*0c10*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 180\n /*0c20*/ LDSM.16.M88.4 R24, [R36+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 182\n /*0c30*/ LDSM.16.MT88.4 R28, [R36+UR4+0x200] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 185\n /*0c40*/ HMMA.16816.F32 R4, R24.reuse, R28, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 188\n /*0c50*/ HMMA.16816.F32 R8, R24, R30, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 193\n /*0c60*/ IMAD.WIDE R24, R37, 0x2, R32 ;\n /*0c70*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 197\n /*0c80*/ @P0 BRA `(.L_x_1) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 338\n\t//## File \".nv_debug_ptx_txt\", line 200\n /*0c90*/ F2FP.F16.F32.PACK_AB R11, R11, R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 201\n /*0ca0*/ F2FP.F16.F32.PACK_AB R10, R9, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 202\n /*0cb0*/ F2FP.F16.F32.PACK_AB R9, R7, R6 ;\n\t//## File \".nv_debug_ptx_txt\", line 203\n /*0cc0*/ F2FP.F16.F32.PACK_AB R8, R5, R4 ;\n.L_x_0:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 228\n /*0cd0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 230\n /*0ce0*/ ISETP.NE.U32.AND P0, PT, R41, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 239\n /*0cf0*/ IMAD.SHL.U32 R3, R0, 0x8, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 242\n /*0d00*/ ISETP.NE.U32.AND P1, PT, R40, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 232\n /*0d10*/ IMAD.SHL.U32 R0, R0, 0x4, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 231\n /*0d20*/ SEL R4, RZ, 0x140, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 224\n /*0d30*/ ULDC.64 UR10, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 240\n /*0d40*/ LOP3.LUT R3, R3, 0xb0, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 216\n /*0d50*/ ULDC UR5, c[0x0][0x23c] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 234\n /*0d60*/ LOP3.LUT R15, R4, 0xb0, R15, 0xf8, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 220\n /*0d70*/ ULDC.64 UR8, c[0x0][0x220] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 224\n /*0d80*/ ISETP.GE.AND P0, PT, R14, UR11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 244\n /*0d90*/ IMAD R7, R2, 0x40, R3 ;\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0da0*/ LOP3.LUT R15, R15, 0x40, R0, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 216\n /*0db0*/ IMAD R0, R13.reuse, UR5, RZ ;\n /*0dc0*/ ISETP.LT.AND P0, PT, R13, UR10, !P0 ;\n\t//## File \".nv_debug_ptx_txt\", line 220\n /*0dd0*/ IMAD.SHL.U32 R3, R14, 0x2, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 243\n /*0de0*/ SEL R4, RZ, 0x140, !P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 218\n /*0df0*/ IMAD.SHL.U32 R2, R0, 0x2, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 236\n /*0e00*/ VIADD R15, R15, UR4 ;\n\t//## File \".nv_debug_ptx_txt\", line 245\n /*0e10*/ LOP3.LUT R4, R7, R4, RZ, 0x3c, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 220\n /*0e20*/ IMAD.HI R5, R14, 0x2, RZ ;\n /*0e30*/ IADD3 R2, P1, P2, R3, UR8, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 237\n /*0e40*/ STSM.16.M88.4 [R15], R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 218\n /*0e50*/ IMAD.HI R0, R0, 0x2, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 220\n /*0e60*/ IADD3.X R3, R5, UR9, R0, P1, P2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 238\n /*0e70*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 249\n /*0e80*/ @!P0 EXIT ;\n\t//## File \".nv_debug_ptx_txt\", line 247\n /*0e90*/ LDS.128 R4, [R4+UR4] ;\n\t//## File \".nv_debug_ptx_txt\", line 249\n /*0ea0*/ STG.E.128 desc[UR6][R2.64], R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 252\n /*0eb0*/ EXIT ;\n.L_x_2:\n /*0ec0*/ BRA `(.L_x_2);\n /*0ed0*/ NOP;\n /*0ee0*/ NOP;\n /*0ef0*/ NOP;\n /*0f00*/ NOP;\n /*0f10*/ NOP;\n /*0f20*/ NOP;\n /*0f30*/ NOP;\n /*0f40*/ NOP;\n /*0f50*/ NOP;\n /*0f60*/ NOP;\n /*0f70*/ NOP;\n.L_x_3:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":266,"end_line":344,"code":" @triton.autotune(\n configs=[\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 16,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 32,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n ],\n key=[\"M\", \"N\", \"K\"],\n )\n @triton.jit\n def matmul_kernel(\n a,\n b,\n c,\n M,\n N,\n K,\n stride_am,\n stride_ak,\n stride_bk,\n stride_bn,\n stride_cm,\n stride_cn,\n BLOCK_SIZE_M: tl.constexpr,\n BLOCK_SIZE_N: tl.constexpr,\n BLOCK_SIZE_K: tl.constexpr,\n GROUP_SIZE_M: tl.constexpr,\n ACTIVATION: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)\n num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)\n num_pid_in_group = GROUP_SIZE_M * num_pid_n\n group_id = pid // num_pid_in_group\n first_pid_m = group_id * GROUP_SIZE_M\n group_size = min(num_pid_m - first_pid_m, GROUP_SIZE_M)\n pid_m = first_pid_m + (pid % group_size)\n pid_n = (pid % num_pid_in_group) // group_size\n\n offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M\n offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N\n offs_k = tl.arange(0, BLOCK_SIZE_K)\n a_ptrs = a + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)\n b_ptrs = b + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)\n\n accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)\n for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):\n a_block = tl.load(\n a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0\n )\n b_block = tl.load(\n b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0\n )\n accumulator += tl.dot(a_block, b_block)\n a_ptrs += BLOCK_SIZE_K * stride_ak\n b_ptrs += BLOCK_SIZE_K * stride_bk\n c_block = accumulator.to(tl.float16)\n\n offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)\n offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)\n c_ptrs = c + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]\n c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)\n tl.store(c_ptrs, c_block, mask=c_mask)\n"}}} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":149,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":160,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":160,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":160,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":160,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":160,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"launch","pid":1724463,"name":"matmul_kernel","function":94769169393936,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"751adc2d5db591acd5d685476d2a0ac4f16c705a3f85b77146db4535cd91888b","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7f16bfe00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":172,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":761,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":328,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ"} -{"event_type":"compilation","pid":1724463,"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":240,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":229,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":164,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":149,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":150,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":739,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":884,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:22.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"f7336a625a0c6f924e5fee4b7063c21567b390bb49fc8cc4fd07cd7d06898f5e","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1536,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[["tt.divisibility",16]],"(5,)":[["tt.divisibility",16]],"(6,)":[["tt.divisibility",16]],"(7,)":[],"(8,)":[["tt.divisibility",16]],"(9,)":[],"(10,)":[["tt.divisibility",16]],"(11,)":[],"(12,)":[],"(13,)":[],"(14,)":[],"(15,)":[],"(16,)":[]},"src_cache_key":"528f8911847b1f39f462157322ec9a674a5565b9b74d621d5ff09625eb710f38","src_constants":{"(7,)":1,"(9,)":1,"(11,)":1,"(12,)":32,"(13,)":16,"(14,)":16,"(15,)":1,"(16,)":null},"times":{"ir_initialization":1150,"lowering_stages":[],"store_results":0}},"file_path":{"matmul_kernel.source":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.source","matmul_kernel.ttir":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.ttir","matmul_kernel.ttgir":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.ttgir","matmul_kernel.llir":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.llir","matmul_kernel.ptx":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.ptx","matmul_kernel.cubin":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.cubin","matmul_kernel.json":"/home/user/.triton/cache/64ZWUYS2BRXZETS75ZFXAY6CCVT3HEF3JH6IZRH5A7GX2BUJR5PA/matmul_kernel.json"},"file_content":{"matmul_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:0)\n#loc56 = loc(\"a\"(#loc))\n#loc57 = loc(\"b\"(#loc))\n#loc58 = loc(\"c\"(#loc))\n#loc59 = loc(\"M\"(#loc))\n#loc60 = loc(\"N\"(#loc))\n#loc61 = loc(\"K\"(#loc))\n#loc62 = loc(\"stride_am\"(#loc))\n#loc63 = loc(\"stride_bk\"(#loc))\n#loc64 = loc(\"stride_cm\"(#loc))\nmodule {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %num_pid_m = arith.constant 31 : i32 loc(#loc115)\n %cst = arith.constant dense<0.000000e+00> : tensor<16x16xf16> loc(#loc1)\n %cst_0 = arith.constant dense<0.000000e+00> : tensor<32x16xf16> loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc3)\n %cst_1 = arith.constant dense<16> : tensor<32x16xi32> loc(#loc1)\n %cst_2 = arith.constant dense<0.000000e+00> : tensor<32x16xf32> loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c32_i32 = arith.constant 32 : i32 loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc66)\n %num_pid_m_3 = arith.addi %M, %num_pid_m : i32 loc(#loc116)\n %num_pid_m_4 = arith.divsi %num_pid_m_3, %c32_i32 : i32 loc(#loc117)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc118)\n %num_pid_n_5 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc119)\n %group_id = arith.divsi %pid, %num_pid_n_5 : i32 loc(#loc68)\n %group_size = arith.subi %num_pid_m_4, %group_id : i32 loc(#loc69)\n %group_size_6 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc70)\n %pid_m = arith.remsi %pid, %group_size_6 : i32 loc(#loc71)\n %pid_m_7 = arith.addi %group_id, %pid_m : i32 loc(#loc72)\n %pid_n = arith.remsi %pid, %num_pid_n_5 : i32 loc(#loc73)\n %pid_n_8 = arith.divsi %pid_n, %group_size_6 : i32 loc(#loc74)\n %offs_am = arith.muli %pid_m_7, %c32_i32 : i32 loc(#loc75)\n %offs_am_9 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc76)\n %offs_am_10 = tt.splat %offs_am : i32 -> tensor<32xi32> loc(#loc77)\n %offs_am_11 = arith.addi %offs_am_10, %offs_am_9 : tensor<32xi32> loc(#loc77)\n %offs_am_12 = tt.splat %M : i32 -> tensor<32xi32> loc(#loc78)\n %offs_am_13 = arith.remsi %offs_am_11, %offs_am_12 : tensor<32xi32> loc(#loc78)\n %offs_bn = arith.muli %pid_n_8, %c16_i32 : i32 loc(#loc79)\n %offs_bn_14 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc80)\n %offs_bn_15 = tt.splat %offs_bn : i32 -> tensor<16xi32> loc(#loc81)\n %offs_bn_16 = arith.addi %offs_bn_15, %offs_bn_14 : tensor<16xi32> loc(#loc81)\n %offs_bn_17 = tt.splat %N : i32 -> tensor<16xi32> loc(#loc82)\n %offs_bn_18 = arith.remsi %offs_bn_16, %offs_bn_17 : tensor<16xi32> loc(#loc82)\n %a_ptrs = tt.expand_dims %offs_am_13 {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc83)\n %a_ptrs_19 = tt.splat %stride_am : i32 -> tensor<32x1xi32> loc(#loc84)\n %a_ptrs_20 = arith.muli %a_ptrs, %a_ptrs_19 : tensor<32x1xi32> loc(#loc84)\n %a_ptrs_21 = tt.expand_dims %offs_bn_14 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc85)\n %a_ptrs_22 = tt.broadcast %a_ptrs_20 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc86)\n %a_ptrs_23 = tt.broadcast %a_ptrs_21 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc86)\n %a_ptrs_24 = arith.addi %a_ptrs_22, %a_ptrs_23 : tensor<32x16xi32> loc(#loc86)\n %a_ptrs_25 = tt.splat %a : !tt.ptr -> tensor<32x16x!tt.ptr> loc(#loc87)\n %a_ptrs_26 = tt.addptr %a_ptrs_25, %a_ptrs_24 : tensor<32x16x!tt.ptr>, tensor<32x16xi32> loc(#loc87)\n %b_ptrs = tt.expand_dims %offs_bn_14 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc88)\n %b_ptrs_27 = tt.splat %stride_bk : i32 -> tensor<16x1xi32> loc(#loc89)\n %b_ptrs_28 = arith.muli %b_ptrs, %b_ptrs_27 : tensor<16x1xi32> loc(#loc89)\n %b_ptrs_29 = tt.expand_dims %offs_bn_18 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc90)\n %b_ptrs_30 = tt.broadcast %b_ptrs_28 : tensor<16x1xi32> -> tensor<16x16xi32> loc(#loc91)\n %b_ptrs_31 = tt.broadcast %b_ptrs_29 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc91)\n %b_ptrs_32 = arith.addi %b_ptrs_30, %b_ptrs_31 : tensor<16x16xi32> loc(#loc91)\n %b_ptrs_33 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr> loc(#loc92)\n %b_ptrs_34 = tt.addptr %b_ptrs_33, %b_ptrs_32 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc92)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc93)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc94)\n %accumulator:3 = scf.for %k = %c0_i32 to %1 step %c1_i32 iter_args(%a_ptrs_49 = %a_ptrs_26, %b_ptrs_50 = %b_ptrs_34, %accumulator_51 = %cst_2) -> (tensor<32x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<32x16xf32>) : i32 {\n %a_block = arith.muli %k, %c16_i32 : i32 loc(#loc96)\n %a_block_52 = arith.subi %K, %a_block : i32 loc(#loc97)\n %a_block_53 = tt.splat %a_block_52 : i32 -> tensor<1x16xi32> loc(#loc98)\n %a_block_54 = arith.cmpi slt, %a_ptrs_21, %a_block_53 : tensor<1x16xi32> loc(#loc98)\n %a_block_55 = tt.broadcast %a_block_54 : tensor<1x16xi1> -> tensor<32x16xi1> loc(#loc99)\n %a_block_56 = tt.load %a_ptrs_49, %a_block_55, %cst_0 : tensor<32x16x!tt.ptr> loc(#loc99)\n %b_block = tt.splat %a_block_52 : i32 -> tensor<16x1xi32> loc(#loc100)\n %b_block_57 = arith.cmpi slt, %b_ptrs, %b_block : tensor<16x1xi32> loc(#loc100)\n %b_block_58 = tt.broadcast %b_block_57 : tensor<16x1xi1> -> tensor<16x16xi1> loc(#loc101)\n %b_block_59 = tt.load %b_ptrs_50, %b_block_58, %cst : tensor<16x16x!tt.ptr> loc(#loc101)\n %accumulator_60 = tt.dot %a_block_56, %b_block_59, %accumulator_51, inputPrecision = tf32 : tensor<32x16xf16> * tensor<16x16xf16> -> tensor<32x16xf32> loc(#loc102)\n %a_ptrs_61 = tt.addptr %a_ptrs_49, %cst_1 : tensor<32x16x!tt.ptr>, tensor<32x16xi32> loc(#loc103)\n %b_ptrs_62 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc104)\n %b_ptrs_63 = tt.splat %b_ptrs_62 : i32 -> tensor<16x16xi32> loc(#loc105)\n %b_ptrs_64 = tt.addptr %b_ptrs_50, %b_ptrs_63 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc105)\n scf.yield %a_ptrs_61, %b_ptrs_64, %accumulator_60 : tensor<32x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<32x16xf32> loc(#loc44)\n } loc(#loc121)\n %c_block = arith.truncf %accumulator#2 : tensor<32x16xf32> to tensor<32x16xf16> loc(#loc106)\n %c_ptrs = tt.expand_dims %offs_am_11 {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc107)\n %c_ptrs_35 = tt.splat %stride_cm : i32 -> tensor<32x1xi32> loc(#loc108)\n %c_ptrs_36 = arith.muli %c_ptrs_35, %c_ptrs : tensor<32x1xi32> loc(#loc108)\n %c_ptrs_37 = tt.splat %c : !tt.ptr -> tensor<32x1x!tt.ptr> loc(#loc109)\n %c_ptrs_38 = tt.addptr %c_ptrs_37, %c_ptrs_36 : tensor<32x1x!tt.ptr>, tensor<32x1xi32> loc(#loc109)\n %c_ptrs_39 = tt.expand_dims %offs_bn_16 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc110)\n %c_ptrs_40 = tt.broadcast %c_ptrs_38 : tensor<32x1x!tt.ptr> -> tensor<32x16x!tt.ptr> loc(#loc111)\n %c_ptrs_41 = tt.broadcast %c_ptrs_39 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc111)\n %c_ptrs_42 = tt.addptr %c_ptrs_40, %c_ptrs_41 : tensor<32x16x!tt.ptr>, tensor<32x16xi32> loc(#loc111)\n %c_mask = tt.splat %M : i32 -> tensor<32x1xi32> loc(#loc112)\n %c_mask_43 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<32x1xi32> loc(#loc112)\n %c_mask_44 = tt.splat %N : i32 -> tensor<1x16xi32> loc(#loc113)\n %c_mask_45 = arith.cmpi slt, %c_ptrs_39, %c_mask_44 : tensor<1x16xi32> loc(#loc113)\n %c_mask_46 = tt.broadcast %c_mask_43 : tensor<32x1xi1> -> tensor<32x16xi1> loc(#loc114)\n %c_mask_47 = tt.broadcast %c_mask_45 : tensor<1x16xi1> -> tensor<32x16xi1> loc(#loc114)\n %c_mask_48 = arith.andi %c_mask_46, %c_mask_47 : tensor<32x16xi1> loc(#loc114)\n tt.store %c_ptrs_42, %c_block, %c_mask_48 : tensor<32x16x!tt.ptr> loc(#loc54)\n tt.return loc(#loc55)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:27)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:22)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:24)\n#loc5 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:17)\n#loc6 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:30)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:27)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:22)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:33)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:46)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:33)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:27)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:19)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:40)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:23)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:51)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:38)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:68)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:23)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:51)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:38)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:68)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:26)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:37)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:56)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:49)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:25)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:36)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:56)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:48)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:33)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:51)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:47)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:43)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:12)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:43)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:12)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:39)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:18)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:33)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:18)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:8)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:29)\n#loc46 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:37)\n#loc47 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:29)\n#loc48 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:17)\n#loc49 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:68)\n#loc50 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:48)\n#loc51 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:33)\n#loc52 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:58)\n#loc53 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:39)\n#loc54 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:21)\n#loc55 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:4)\n#loc65 = loc(\"num_pid_m\"(#loc2))\n#loc66 = loc(\"pid\"(#loc4))\n#loc67 = loc(\"num_pid_n\"(#loc7))\n#loc68 = loc(\"group_id\"(#loc8))\n#loc69 = loc(\"group_size\"(#loc9))\n#loc70 = loc(\"group_size\"(#loc10))\n#loc71 = loc(\"pid_m\"(#loc11))\n#loc72 = loc(\"pid_m\"(#loc12))\n#loc73 = loc(\"pid_n\"(#loc13))\n#loc74 = loc(\"pid_n\"(#loc14))\n#loc75 = loc(\"offs_am\"(#loc15))\n#loc76 = loc(\"offs_am\"(#loc16))\n#loc77 = loc(\"offs_am\"(#loc17))\n#loc78 = loc(\"offs_am\"(#loc18))\n#loc79 = loc(\"offs_bn\"(#loc19))\n#loc80 = loc(\"offs_bn\"(#loc20))\n#loc81 = loc(\"offs_bn\"(#loc21))\n#loc82 = loc(\"offs_bn\"(#loc22))\n#loc83 = loc(\"a_ptrs\"(#loc23))\n#loc84 = loc(\"a_ptrs\"(#loc24))\n#loc85 = loc(\"a_ptrs\"(#loc25))\n#loc86 = loc(\"a_ptrs\"(#loc26))\n#loc87 = loc(\"a_ptrs\"(#loc27))\n#loc88 = loc(\"b_ptrs\"(#loc28))\n#loc89 = loc(\"b_ptrs\"(#loc29))\n#loc90 = loc(\"b_ptrs\"(#loc30))\n#loc91 = loc(\"b_ptrs\"(#loc31))\n#loc92 = loc(\"b_ptrs\"(#loc32))\n#loc93 = loc(callsite(#loc5 at #loc33))\n#loc94 = loc(callsite(#loc6 at #loc33))\n#loc95 = loc(\"a_ptrs\"(#loc3))\n#loc96 = loc(\"a_block\"(#loc34))\n#loc97 = loc(\"a_block\"(#loc35))\n#loc98 = loc(\"a_block\"(#loc36))\n#loc99 = loc(\"a_block\"(#loc37))\n#loc100 = loc(\"b_block\"(#loc38))\n#loc101 = loc(\"b_block\"(#loc39))\n#loc102 = loc(\"accumulator\"(#loc40))\n#loc103 = loc(\"a_ptrs\"(#loc41))\n#loc104 = loc(\"b_ptrs\"(#loc42))\n#loc105 = loc(\"b_ptrs\"(#loc43))\n#loc106 = loc(\"c_block\"(#loc45))\n#loc107 = loc(\"c_ptrs\"(#loc46))\n#loc108 = loc(\"c_ptrs\"(#loc47))\n#loc109 = loc(\"c_ptrs\"(#loc48))\n#loc110 = loc(\"c_ptrs\"(#loc49))\n#loc111 = loc(\"c_ptrs\"(#loc50))\n#loc112 = loc(\"c_mask\"(#loc51))\n#loc113 = loc(\"c_mask\"(#loc52))\n#loc114 = loc(\"c_mask\"(#loc53))\n#loc115 = loc(callsite(#loc1 at #loc65))\n#loc116 = loc(callsite(#loc5 at #loc65))\n#loc117 = loc(callsite(#loc6 at #loc65))\n#loc118 = loc(callsite(#loc5 at #loc67))\n#loc119 = loc(callsite(#loc6 at #loc67))\n#loc120 = loc(\"b_ptrs\"(#loc95))\n#loc121 = loc(\"accumulator\"(#loc120))\n","matmul_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1, 8], threadsPerWarp = [16, 2], warpsPerCTA = [1, 1], order = [1, 0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:0)\n#mma = #ttg.nvidia_mma<{versionMajor = 2, versionMinor = 0, warpsPerCTA = [1, 1], instrShape = [16, 8]}>\n#shared = #ttg.swizzled_shared<{vec = 8, perPhase = 4, maxPhase = 2, order = [1, 0]}>\n#smem = #ttg.shared_memory\n#loc56 = loc(\"a\"(#loc))\n#loc57 = loc(\"b\"(#loc))\n#loc58 = loc(\"c\"(#loc))\n#loc59 = loc(\"M\"(#loc))\n#loc60 = loc(\"N\"(#loc))\n#loc61 = loc(\"K\"(#loc))\n#loc62 = loc(\"stride_am\"(#loc))\n#loc63 = loc(\"stride_bk\"(#loc))\n#loc64 = loc(\"stride_cm\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 1 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %cst = arith.constant dense<0.000000e+00> : tensor<32x16xf32, #mma> loc(#loc1)\n %cst_0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #blocked> loc(#loc1)\n %cst_1 = arith.constant dense<0.000000e+00> : tensor<32x16xf16, #blocked> loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %c32_i32 = arith.constant 32 : i32 loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %c31_i32 = arith.constant 31 : i32 loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc1)\n %cst_2 = arith.constant dense<16> : tensor<32x16xi32, #blocked> loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc65)\n %num_pid_m = arith.addi %M, %c31_i32 : i32 loc(#loc115)\n %num_pid_m_3 = arith.divsi %num_pid_m, %c32_i32 : i32 loc(#loc116)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc117)\n %num_pid_n_4 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc118)\n %group_id = arith.divsi %pid, %num_pid_n_4 : i32 loc(#loc68)\n %group_size = arith.subi %num_pid_m_3, %group_id : i32 loc(#loc69)\n %group_size_5 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc70)\n %pid_m = arith.remsi %pid, %group_size_5 : i32 loc(#loc71)\n %pid_m_6 = arith.addi %group_id, %pid_m : i32 loc(#loc72)\n %pid_n = arith.remsi %pid, %num_pid_n_4 : i32 loc(#loc73)\n %pid_n_7 = arith.divsi %pid_n, %group_size_5 : i32 loc(#loc74)\n %offs_am = arith.muli %pid_m_6, %c32_i32 : i32 loc(#loc75)\n %offs_am_8 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc76)\n %offs_am_9 = tt.splat %offs_am : i32 -> tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc77)\n %offs_am_10 = arith.addi %offs_am_9, %offs_am_8 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc77)\n %offs_am_11 = tt.splat %M : i32 -> tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc78)\n %offs_am_12 = arith.remsi %offs_am_10, %offs_am_11 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc78)\n %offs_bn = arith.muli %pid_n_7, %c16_i32 : i32 loc(#loc79)\n %offs_bn_13 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc80)\n %offs_bn_14 = tt.splat %offs_bn : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc81)\n %offs_bn_15 = arith.addi %offs_bn_14, %offs_bn_13 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc81)\n %offs_bn_16 = tt.splat %N : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc82)\n %offs_bn_17 = arith.remsi %offs_bn_15, %offs_bn_16 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc82)\n %a_ptrs = tt.expand_dims %offs_am_12 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked> loc(#loc83)\n %a_ptrs_18 = tt.splat %stride_am : i32 -> tensor<32x1xi32, #blocked> loc(#loc84)\n %a_ptrs_19 = arith.muli %a_ptrs, %a_ptrs_18 : tensor<32x1xi32, #blocked> loc(#loc84)\n %a_ptrs_20 = tt.expand_dims %offs_bn_13 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc85)\n %a_ptrs_21 = tt.broadcast %a_ptrs_19 : tensor<32x1xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc86)\n %a_ptrs_22 = tt.broadcast %a_ptrs_20 : tensor<1x16xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc86)\n %a_ptrs_23 = arith.addi %a_ptrs_21, %a_ptrs_22 : tensor<32x16xi32, #blocked> loc(#loc86)\n %a_ptrs_24 = tt.splat %a : !tt.ptr -> tensor<32x16x!tt.ptr, #blocked> loc(#loc87)\n %a_ptrs_25 = tt.addptr %a_ptrs_24, %a_ptrs_23 : tensor<32x16x!tt.ptr, #blocked>, tensor<32x16xi32, #blocked> loc(#loc87)\n %b_ptrs = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc88)\n %b_ptrs_26 = tt.expand_dims %b_ptrs {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc88)\n %b_ptrs_27 = tt.splat %stride_bk : i32 -> tensor<16x1xi32, #blocked> loc(#loc89)\n %b_ptrs_28 = arith.muli %b_ptrs_26, %b_ptrs_27 : tensor<16x1xi32, #blocked> loc(#loc89)\n %b_ptrs_29 = tt.expand_dims %offs_bn_17 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc90)\n %b_ptrs_30 = tt.broadcast %b_ptrs_28 : tensor<16x1xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc91)\n %b_ptrs_31 = tt.broadcast %b_ptrs_29 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc91)\n %b_ptrs_32 = arith.addi %b_ptrs_30, %b_ptrs_31 : tensor<16x16xi32, #blocked> loc(#loc91)\n %b_ptrs_33 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr, #blocked> loc(#loc92)\n %b_ptrs_34 = tt.addptr %b_ptrs_33, %b_ptrs_32 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc92)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc93)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc94)\n %b_ptrs_35 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc95)\n %b_ptrs_36 = tt.splat %b_ptrs_35 : i32 -> tensor<16x16xi32, #blocked> loc(#loc96)\n %accumulator:3 = scf.for %accumulator_51 = %c0_i32 to %1 step %c1_i32 iter_args(%arg10 = %cst, %a_ptrs_52 = %a_ptrs_25, %b_ptrs_53 = %b_ptrs_34) -> (tensor<32x16xf32, #mma>, tensor<32x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked>) : i32 {\n %a_block = arith.muli %accumulator_51, %c16_i32 : i32 loc(#loc98)\n %a_block_54 = arith.subi %K, %a_block : i32 loc(#loc99)\n %a_block_55 = tt.splat %a_block_54 : i32 -> tensor<1x16xi32, #blocked> loc(#loc100)\n %a_block_56 = arith.cmpi slt, %a_ptrs_20, %a_block_55 : tensor<1x16xi32, #blocked> loc(#loc100)\n %a_block_57 = tt.broadcast %a_block_56 : tensor<1x16xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc101)\n %a_block_58 = tt.load %a_ptrs_52, %a_block_57, %cst_1 : tensor<32x16x!tt.ptr, #blocked> loc(#loc101)\n %a_block_59 = ttg.local_alloc %a_block_58 : (tensor<32x16xf16, #blocked>) -> !ttg.memdesc<32x16xf16, #shared, #smem> loc(#loc101)\n %b_block = tt.splat %a_block_54 : i32 -> tensor<16x1xi32, #blocked> loc(#loc102)\n %b_block_60 = arith.cmpi slt, %b_ptrs_26, %b_block : tensor<16x1xi32, #blocked> loc(#loc102)\n %b_block_61 = tt.broadcast %b_block_60 : tensor<16x1xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc103)\n %b_block_62 = tt.load %b_ptrs_53, %b_block_61, %cst_0 : tensor<16x16x!tt.ptr, #blocked> loc(#loc103)\n %b_block_63 = ttg.local_alloc %b_block_62 : (tensor<16x16xf16, #blocked>) -> !ttg.memdesc<16x16xf16, #shared, #smem> loc(#loc103)\n %a_block_64 = ttg.local_load %a_block_59 : !ttg.memdesc<32x16xf16, #shared, #smem> -> tensor<32x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> loc(#loc101)\n %b_block_65 = ttg.local_load %b_block_63 : !ttg.memdesc<16x16xf16, #shared, #smem> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> loc(#loc103)\n %accumulator_66 = tt.dot %a_block_64, %b_block_65, %arg10, inputPrecision = tf32 : tensor<32x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> * tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> -> tensor<32x16xf32, #mma> loc(#loc104)\n %a_ptrs_67 = tt.addptr %a_ptrs_52, %cst_2 : tensor<32x16x!tt.ptr, #blocked>, tensor<32x16xi32, #blocked> loc(#loc105)\n %b_ptrs_68 = tt.addptr %b_ptrs_53, %b_ptrs_36 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc96)\n scf.yield %accumulator_66, %a_ptrs_67, %b_ptrs_68 : tensor<32x16xf32, #mma>, tensor<32x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked> loc(#loc44)\n } loc(#loc120)\n %c_block = arith.truncf %accumulator#0 : tensor<32x16xf32, #mma> to tensor<32x16xf16, #mma> loc(#loc106)\n %c_ptrs = tt.expand_dims %offs_am_10 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked> loc(#loc107)\n %c_ptrs_37 = tt.splat %stride_cm : i32 -> tensor<32x1xi32, #blocked> loc(#loc108)\n %c_ptrs_38 = arith.muli %c_ptrs_37, %c_ptrs : tensor<32x1xi32, #blocked> loc(#loc108)\n %c_ptrs_39 = tt.splat %c : !tt.ptr -> tensor<32x1x!tt.ptr, #blocked> loc(#loc109)\n %c_ptrs_40 = tt.addptr %c_ptrs_39, %c_ptrs_38 : tensor<32x1x!tt.ptr, #blocked>, tensor<32x1xi32, #blocked> loc(#loc109)\n %c_ptrs_41 = tt.expand_dims %offs_bn_15 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc110)\n %c_ptrs_42 = tt.broadcast %c_ptrs_40 : tensor<32x1x!tt.ptr, #blocked> -> tensor<32x16x!tt.ptr, #blocked> loc(#loc111)\n %c_ptrs_43 = tt.broadcast %c_ptrs_41 : tensor<1x16xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc111)\n %c_ptrs_44 = tt.addptr %c_ptrs_42, %c_ptrs_43 : tensor<32x16x!tt.ptr, #blocked>, tensor<32x16xi32, #blocked> loc(#loc111)\n %c_mask = tt.splat %M : i32 -> tensor<32x1xi32, #blocked> loc(#loc112)\n %c_mask_45 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<32x1xi32, #blocked> loc(#loc112)\n %c_mask_46 = tt.splat %N : i32 -> tensor<1x16xi32, #blocked> loc(#loc113)\n %c_mask_47 = arith.cmpi slt, %c_ptrs_41, %c_mask_46 : tensor<1x16xi32, #blocked> loc(#loc113)\n %c_mask_48 = tt.broadcast %c_mask_45 : tensor<32x1xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc114)\n %c_mask_49 = tt.broadcast %c_mask_47 : tensor<1x16xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc114)\n %c_mask_50 = arith.andi %c_mask_48, %c_mask_49 : tensor<32x16xi1, #blocked> loc(#loc114)\n %2 = ttg.convert_layout %c_block : tensor<32x16xf16, #mma> -> tensor<32x16xf16, #blocked> loc(#loc54)\n tt.store %c_ptrs_44, %2, %c_mask_50 : tensor<32x16x!tt.ptr, #blocked> loc(#loc54)\n tt.return loc(#loc55)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:24)\n#loc3 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:17)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:27)\n#loc5 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:30)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:27)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:22)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:33)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:46)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:33)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:27)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:19)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:40)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:23)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:51)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:38)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:68)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:23)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:51)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:38)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:68)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:26)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:37)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:56)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:49)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:25)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:36)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:56)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:48)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:33)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:33)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:18)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:22)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:51)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:47)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:43)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:12)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:43)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:12)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:39)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:18)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:8)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:29)\n#loc46 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:37)\n#loc47 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:29)\n#loc48 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:17)\n#loc49 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:68)\n#loc50 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:48)\n#loc51 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:33)\n#loc52 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:58)\n#loc53 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:39)\n#loc54 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:21)\n#loc55 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:4)\n#loc65 = loc(\"pid\"(#loc2))\n#loc66 = loc(\"num_pid_m\"(#loc4))\n#loc67 = loc(\"num_pid_n\"(#loc6))\n#loc68 = loc(\"group_id\"(#loc7))\n#loc69 = loc(\"group_size\"(#loc8))\n#loc70 = loc(\"group_size\"(#loc9))\n#loc71 = loc(\"pid_m\"(#loc10))\n#loc72 = loc(\"pid_m\"(#loc11))\n#loc73 = loc(\"pid_n\"(#loc12))\n#loc74 = loc(\"pid_n\"(#loc13))\n#loc75 = loc(\"offs_am\"(#loc14))\n#loc76 = loc(\"offs_am\"(#loc15))\n#loc77 = loc(\"offs_am\"(#loc16))\n#loc78 = loc(\"offs_am\"(#loc17))\n#loc79 = loc(\"offs_bn\"(#loc18))\n#loc80 = loc(\"offs_bn\"(#loc19))\n#loc81 = loc(\"offs_bn\"(#loc20))\n#loc82 = loc(\"offs_bn\"(#loc21))\n#loc83 = loc(\"a_ptrs\"(#loc22))\n#loc84 = loc(\"a_ptrs\"(#loc23))\n#loc85 = loc(\"a_ptrs\"(#loc24))\n#loc86 = loc(\"a_ptrs\"(#loc25))\n#loc87 = loc(\"a_ptrs\"(#loc26))\n#loc88 = loc(\"b_ptrs\"(#loc27))\n#loc89 = loc(\"b_ptrs\"(#loc28))\n#loc90 = loc(\"b_ptrs\"(#loc29))\n#loc91 = loc(\"b_ptrs\"(#loc30))\n#loc92 = loc(\"b_ptrs\"(#loc31))\n#loc93 = loc(callsite(#loc3 at #loc32))\n#loc94 = loc(callsite(#loc5 at #loc32))\n#loc95 = loc(\"b_ptrs\"(#loc33))\n#loc96 = loc(\"b_ptrs\"(#loc34))\n#loc97 = loc(\"a_ptrs\"(#loc35))\n#loc98 = loc(\"a_block\"(#loc36))\n#loc99 = loc(\"a_block\"(#loc37))\n#loc100 = loc(\"a_block\"(#loc38))\n#loc101 = loc(\"a_block\"(#loc39))\n#loc102 = loc(\"b_block\"(#loc40))\n#loc103 = loc(\"b_block\"(#loc41))\n#loc104 = loc(\"accumulator\"(#loc42))\n#loc105 = loc(\"a_ptrs\"(#loc43))\n#loc106 = loc(\"c_block\"(#loc45))\n#loc107 = loc(\"c_ptrs\"(#loc46))\n#loc108 = loc(\"c_ptrs\"(#loc47))\n#loc109 = loc(\"c_ptrs\"(#loc48))\n#loc110 = loc(\"c_ptrs\"(#loc49))\n#loc111 = loc(\"c_ptrs\"(#loc50))\n#loc112 = loc(\"c_mask\"(#loc51))\n#loc113 = loc(\"c_mask\"(#loc52))\n#loc114 = loc(\"c_mask\"(#loc53))\n#loc115 = loc(callsite(#loc3 at #loc66))\n#loc116 = loc(callsite(#loc5 at #loc66))\n#loc117 = loc(callsite(#loc3 at #loc67))\n#loc118 = loc(callsite(#loc5 at #loc67))\n#loc119 = loc(\"b_ptrs\"(#loc97))\n#loc120 = loc(\"accumulator\"(#loc119))\n","matmul_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n@global_smem = external addrspace(3) global [0 x i8], align 16\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @matmul_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i32 %8, ptr addrspace(1) readnone captures(none) %9, ptr addrspace(1) readnone captures(none) %10) local_unnamed_addr #0 !dbg !4 {\n %12 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %13 = add i32 %3, 31, !dbg !13\n %14 = sdiv i32 %13, 32, !dbg !18\n %15 = add i32 %4, 15, !dbg !19\n %16 = sdiv i32 %15, 16, !dbg !21\n %.frozen = freeze i32 %16, !dbg !22\n %17 = sdiv i32 %12, %.frozen, !dbg !22\n %18 = sub i32 %14, %17, !dbg !23\n %19 = tail call i32 @llvm.smin.i32(i32 %18, i32 1), !dbg !24\n %20 = srem i32 %12, %19, !dbg !25\n %21 = add i32 %20, %17, !dbg !26\n %22 = mul i32 %17, %.frozen, !dbg !27\n %.decomposed = sub i32 %12, %22, !dbg !27\n %23 = sdiv i32 %.decomposed, %19, !dbg !28\n %24 = shl i32 %21, 5, !dbg !29\n %25 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !30\n %26 = lshr i32 %25, 1, !dbg !30\n %27 = and i32 %26, 15, !dbg !30\n %28 = or disjoint i32 %24, %27, !dbg !31\n %29 = or disjoint i32 %28, 16, !dbg !31\n %30 = shl nsw i32 %23, 4, !dbg !32\n %31 = and i32 %25, 1, !dbg !33\n %32 = shl nuw nsw i32 %31, 3, !dbg !33\n %33 = or disjoint i32 %30, %32, !dbg !34\n %34 = add i32 %5, 15, !dbg !35\n %35 = sdiv i32 %34, 16, !dbg !37\n %36 = icmp sgt i32 %34, 15, !dbg !38\n br i1 %36, label %.lr.ph, label %.._crit_edge_crit_edge, !dbg !38\n\n.._crit_edge_crit_edge: ; preds = %11\n %.pre = shl nuw nsw i32 %25, 4, !dbg !39\n %.pre109 = and i32 %25, 4, !dbg !39\n %.pre111 = and i32 %25, 8, !dbg !39\n br label %._crit_edge, !dbg !38\n\n.lr.ph: ; preds = %11\n %37 = shl i32 %7, 4, !dbg !40\n %38 = srem i32 %33, %4, !dbg !41\n %39 = mul i32 %7, %27, !dbg !42\n %40 = add i32 %38, %39, !dbg !43\n %41 = sext i32 %40 to i64, !dbg !44\n %42 = getelementptr half, ptr addrspace(1) %1, i64 %41, !dbg !44\n %43 = srem i32 %29, %3, !dbg !45\n %44 = mul i32 %43, %6, !dbg !46\n %45 = add i32 %44, %32, !dbg !47\n %46 = sext i32 %45 to i64, !dbg !48\n %47 = getelementptr half, ptr addrspace(1) %0, i64 %46, !dbg !48\n %48 = srem i32 %28, %3, !dbg !45\n %49 = mul i32 %48, %6, !dbg !46\n %50 = add i32 %49, %32, !dbg !47\n %51 = sext i32 %50 to i64, !dbg !48\n %52 = getelementptr half, ptr addrspace(1) %0, i64 %51, !dbg !48\n %53 = shl nuw nsw i32 %25, 4\n %54 = and i32 %53, 368\n %55 = and i32 %25, 8\n %56 = icmp eq i32 %55, 0\n %57 = select i1 %56, i32 0, i32 144\n %58 = xor i32 %57, %54\n %59 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %58\n %60 = getelementptr inbounds nuw i8, ptr addrspace(3) %59, i32 512\n %61 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 1024), i32 %58\n %62 = shl nuw nsw i32 %25, 5\n %63 = and i32 %62, 352\n %64 = and i32 %25, 4\n %65 = icmp eq i32 %64, 0\n %66 = select i1 %65, i32 0, i32 144\n %67 = and i32 %25, 16\n %68 = or disjoint i32 %66, %63\n %69 = xor i32 %68, %67\n %70 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %69\n %71 = getelementptr inbounds nuw i8, ptr addrspace(3) %70, i32 512\n %72 = xor i32 %66, %67\n %73 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 1024), i32 %63\n %74 = getelementptr inbounds nuw i8, ptr addrspace(3) %73, i32 %72\n %75 = sext i32 %37 to i64\n br label %76, !dbg !38\n\n76: ; preds = %.lr.ph, %76\n %.pn47107 = phi ptr addrspace(1) [ %42, %.lr.ph ], [ %150, %76 ]\n %.pn15106 = phi ptr addrspace(1) [ %47, %.lr.ph ], [ %149, %76 ]\n %.pn31105 = phi ptr addrspace(1) [ %52, %.lr.ph ], [ %148, %76 ]\n %77 = phi float [ 0.000000e+00, %.lr.ph ], [ %129, %76 ]\n %78 = phi float [ 0.000000e+00, %.lr.ph ], [ %130, %76 ]\n %79 = phi float [ 0.000000e+00, %.lr.ph ], [ %131, %76 ]\n %80 = phi float [ 0.000000e+00, %.lr.ph ], [ %132, %76 ]\n %81 = phi float [ 0.000000e+00, %.lr.ph ], [ %134, %76 ]\n %82 = phi float [ 0.000000e+00, %.lr.ph ], [ %135, %76 ]\n %83 = phi float [ 0.000000e+00, %.lr.ph ], [ %136, %76 ]\n %84 = phi float [ 0.000000e+00, %.lr.ph ], [ %137, %76 ]\n %85 = phi float [ 0.000000e+00, %.lr.ph ], [ %139, %76 ]\n %86 = phi float [ 0.000000e+00, %.lr.ph ], [ %140, %76 ]\n %87 = phi float [ 0.000000e+00, %.lr.ph ], [ %141, %76 ]\n %88 = phi float [ 0.000000e+00, %.lr.ph ], [ %142, %76 ]\n %89 = phi float [ 0.000000e+00, %.lr.ph ], [ %144, %76 ]\n %90 = phi float [ 0.000000e+00, %.lr.ph ], [ %145, %76 ]\n %91 = phi float [ 0.000000e+00, %.lr.ph ], [ %146, %76 ]\n %92 = phi float [ 0.000000e+00, %.lr.ph ], [ %147, %76 ]\n %93 = phi i32 [ 0, %.lr.ph ], [ %151, %76 ]\n %94 = shl i32 %93, 4, !dbg !49\n %95 = sub i32 %5, %94, !dbg !50\n %96 = icmp slt i32 %32, %95, !dbg !51\n %97 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn31105, i1 %96) #6, !dbg !52\n %98 = extractvalue { i32, i32, i32, i32 } %97, 0, !dbg !52\n %99 = extractvalue { i32, i32, i32, i32 } %97, 1, !dbg !52\n %100 = extractvalue { i32, i32, i32, i32 } %97, 2, !dbg !52\n %101 = extractvalue { i32, i32, i32, i32 } %97, 3, !dbg !52\n %102 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn15106, i1 %96) #6, !dbg !52\n %103 = extractvalue { i32, i32, i32, i32 } %102, 0, !dbg !52\n %104 = extractvalue { i32, i32, i32, i32 } %102, 1, !dbg !52\n %105 = extractvalue { i32, i32, i32, i32 } %102, 2, !dbg !52\n %106 = extractvalue { i32, i32, i32, i32 } %102, 3, !dbg !52\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !52\n tail call void asm sideeffect \"st.shared.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %59, i32 %98, i32 %99, i32 %100, i32 %101) #6, !dbg !52\n tail call void asm sideeffect \"st.shared.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) nonnull %60, i32 %103, i32 %104, i32 %105, i32 %106) #6, !dbg !52\n %107 = icmp slt i32 %27, %95, !dbg !53\n %108 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn47107, i1 %107) #6, !dbg !54\n %109 = extractvalue { i32, i32, i32, i32 } %108, 0, !dbg !54\n %110 = extractvalue { i32, i32, i32, i32 } %108, 1, !dbg !54\n %111 = extractvalue { i32, i32, i32, i32 } %108, 2, !dbg !54\n %112 = extractvalue { i32, i32, i32, i32 } %108, 3, !dbg !54\n tail call void asm sideeffect \"st.shared.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %61, i32 %109, i32 %110, i32 %111, i32 %112) #6, !dbg !54\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !52\n %113 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %70), !dbg !52\n %114 = extractvalue { i32, i32, i32, i32 } %113, 0, !dbg !52\n %115 = extractvalue { i32, i32, i32, i32 } %113, 1, !dbg !52\n %116 = extractvalue { i32, i32, i32, i32 } %113, 2, !dbg !52\n %117 = extractvalue { i32, i32, i32, i32 } %113, 3, !dbg !52\n %118 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) nonnull %71), !dbg !52\n %119 = extractvalue { i32, i32, i32, i32 } %118, 0, !dbg !52\n %120 = extractvalue { i32, i32, i32, i32 } %118, 1, !dbg !52\n %121 = extractvalue { i32, i32, i32, i32 } %118, 2, !dbg !52\n %122 = extractvalue { i32, i32, i32, i32 } %118, 3, !dbg !52\n %123 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %74), !dbg !54\n %124 = extractvalue { i32, i32, i32, i32 } %123, 0, !dbg !54\n %125 = extractvalue { i32, i32, i32, i32 } %123, 1, !dbg !54\n %126 = extractvalue { i32, i32, i32, i32 } %123, 2, !dbg !54\n %127 = extractvalue { i32, i32, i32, i32 } %123, 3, !dbg !54\n %128 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %77, float %78, float %79, float %80, i32 %114, i32 %115, i32 %116, i32 %117, i32 %124, i32 %125) #6, !dbg !55\n %129 = extractvalue { float, float, float, float } %128, 0, !dbg !55\n %130 = extractvalue { float, float, float, float } %128, 1, !dbg !55\n %131 = extractvalue { float, float, float, float } %128, 2, !dbg !55\n %132 = extractvalue { float, float, float, float } %128, 3, !dbg !55\n %133 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %81, float %82, float %83, float %84, i32 %114, i32 %115, i32 %116, i32 %117, i32 %126, i32 %127) #6, !dbg !55\n %134 = extractvalue { float, float, float, float } %133, 0, !dbg !55\n %135 = extractvalue { float, float, float, float } %133, 1, !dbg !55\n %136 = extractvalue { float, float, float, float } %133, 2, !dbg !55\n %137 = extractvalue { float, float, float, float } %133, 3, !dbg !55\n %138 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %85, float %86, float %87, float %88, i32 %119, i32 %120, i32 %121, i32 %122, i32 %124, i32 %125) #6, !dbg !55\n %139 = extractvalue { float, float, float, float } %138, 0, !dbg !55\n %140 = extractvalue { float, float, float, float } %138, 1, !dbg !55\n %141 = extractvalue { float, float, float, float } %138, 2, !dbg !55\n %142 = extractvalue { float, float, float, float } %138, 3, !dbg !55\n %143 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %89, float %90, float %91, float %92, i32 %119, i32 %120, i32 %121, i32 %122, i32 %126, i32 %127) #6, !dbg !55\n %144 = extractvalue { float, float, float, float } %143, 0, !dbg !55\n %145 = extractvalue { float, float, float, float } %143, 1, !dbg !55\n %146 = extractvalue { float, float, float, float } %143, 2, !dbg !55\n %147 = extractvalue { float, float, float, float } %143, 3, !dbg !55\n %148 = getelementptr i8, ptr addrspace(1) %.pn31105, i64 32, !dbg !56\n %149 = getelementptr i8, ptr addrspace(1) %.pn15106, i64 32, !dbg !56\n %150 = getelementptr half, ptr addrspace(1) %.pn47107, i64 %75, !dbg !57\n %151 = add nuw nsw i32 %93, 1, !dbg !38\n %exitcond.not = icmp eq i32 %151, %35, !dbg !38\n br i1 %exitcond.not, label %._crit_edge.loopexit, label %76, !dbg !38\n\n._crit_edge.loopexit: ; preds = %76\n %152 = insertelement <16 x float> poison, float %129, i64 0, !dbg !58\n %153 = insertelement <16 x float> %152, float %130, i64 1, !dbg !58\n %154 = insertelement <16 x float> %153, float %131, i64 2, !dbg !58\n %155 = insertelement <16 x float> %154, float %132, i64 3, !dbg !58\n %156 = insertelement <16 x float> %155, float %134, i64 4, !dbg !58\n %157 = insertelement <16 x float> %156, float %135, i64 5, !dbg !58\n %158 = insertelement <16 x float> %157, float %136, i64 6, !dbg !58\n %159 = insertelement <16 x float> %158, float %137, i64 7, !dbg !58\n %160 = insertelement <16 x float> %159, float %139, i64 8, !dbg !58\n %161 = insertelement <16 x float> %160, float %140, i64 9, !dbg !58\n %162 = insertelement <16 x float> %161, float %141, i64 10, !dbg !58\n %163 = insertelement <16 x float> %162, float %142, i64 11, !dbg !58\n %164 = insertelement <16 x float> %163, float %144, i64 12, !dbg !58\n %165 = insertelement <16 x float> %164, float %145, i64 13, !dbg !58\n %166 = insertelement <16 x float> %165, float %146, i64 14, !dbg !58\n %167 = insertelement <16 x float> %166, float %147, i64 15, !dbg !58\n %168 = fptrunc <16 x float> %167 to <16 x half>, !dbg !58\n %169 = bitcast <16 x half> %168 to <8 x i32>, !dbg !39\n br label %._crit_edge, !dbg !58\n\n._crit_edge: ; preds = %._crit_edge.loopexit, %.._crit_edge_crit_edge\n %.pre-phi112 = phi i32 [ %.pre111, %.._crit_edge_crit_edge ], [ %55, %._crit_edge.loopexit ], !dbg !39\n %.pre-phi110 = phi i32 [ %.pre109, %.._crit_edge_crit_edge ], [ %64, %._crit_edge.loopexit ], !dbg !39\n %.pre-phi = phi i32 [ %.pre, %.._crit_edge_crit_edge ], [ %53, %._crit_edge.loopexit ], !dbg !39\n %.bc142 = phi <8 x i32> [ zeroinitializer, %.._crit_edge_crit_edge ], [ %169, %._crit_edge.loopexit ]\n %170 = mul i32 %28, %8, !dbg !59\n %171 = mul i32 %8, %29, !dbg !59\n %172 = sext i32 %170 to i64, !dbg !60\n %173 = getelementptr half, ptr addrspace(1) %2, i64 %172, !dbg !60\n %174 = sext i32 %171 to i64, !dbg !60\n %175 = getelementptr half, ptr addrspace(1) %2, i64 %174, !dbg !60\n %176 = sext i32 %33 to i64, !dbg !61\n %177 = getelementptr half, ptr addrspace(1) %173, i64 %176, !dbg !61\n %178 = getelementptr half, ptr addrspace(1) %175, i64 %176, !dbg !61\n %179 = icmp slt i32 %28, %3, !dbg !62\n %180 = icmp slt i32 %29, %3, !dbg !62\n %181 = icmp slt i32 %33, %4, !dbg !63\n %182 = and i1 %179, %181, !dbg !64\n %183 = and i1 %180, %181, !dbg !64\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !39\n %184 = and i32 %.pre-phi, 176, !dbg !39\n %185 = icmp eq i32 %.pre-phi110, 0, !dbg !39\n %186 = select i1 %185, i32 0, i32 320, !dbg !39\n %187 = shl nuw nsw i32 %25, 2, !dbg !39\n %188 = and i32 %187, 64, !dbg !39\n %189 = or disjoint i32 %186, %184, !dbg !39\n %190 = xor i32 %189, %188, !dbg !39\n %191 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %190, !dbg !39\n %.extract129 = extractelement <8 x i32> %.bc142, i64 0, !dbg !39\n %.extract131 = extractelement <8 x i32> %.bc142, i64 1, !dbg !39\n %.extract133 = extractelement <8 x i32> %.bc142, i64 2, !dbg !39\n %.extract135 = extractelement <8 x i32> %.bc142, i64 3, !dbg !39\n tail call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %191, i32 %.extract129, i32 %.extract131, i32 %.extract133, i32 %.extract135), !dbg !39\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !39\n %192 = shl nuw nsw i32 %25, 3, !dbg !39\n %193 = and i32 %192, 176, !dbg !39\n %194 = shl nuw nsw i32 %31, 6, !dbg !39\n %195 = icmp eq i32 %.pre-phi112, 0, !dbg !39\n %196 = select i1 %195, i32 0, i32 320, !dbg !39\n %197 = or disjoint i32 %193, %194, !dbg !39\n %198 = xor i32 %197, %196, !dbg !39\n %199 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %198, !dbg !39\n %200 = load <4 x i32>, ptr addrspace(3) %199, align 16, !dbg !39\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !39\n %.extract137 = extractelement <8 x i32> %.bc142, i64 4, !dbg !39\n %.extract139 = extractelement <8 x i32> %.bc142, i64 5, !dbg !39\n %.extract141 = extractelement <8 x i32> %.bc142, i64 6, !dbg !39\n %.extract143 = extractelement <8 x i32> %.bc142, i64 7, !dbg !39\n tail call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %191, i32 %.extract137, i32 %.extract139, i32 %.extract141, i32 %.extract143), !dbg !39\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !39\n %201 = load <4 x i32>, ptr addrspace(3) %199, align 16, !dbg !39\n %.extract = extractelement <4 x i32> %200, i64 0, !dbg !39\n %.extract92 = extractelement <4 x i32> %200, i64 1, !dbg !39\n %.extract94 = extractelement <4 x i32> %200, i64 2, !dbg !39\n %.extract96 = extractelement <4 x i32> %200, i64 3, !dbg !39\n tail call void asm sideeffect \"@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };\", \"r,r,r,r,l,b\"(i32 %.extract, i32 %.extract92, i32 %.extract94, i32 %.extract96, ptr addrspace(1) %177, i1 %182) #6, !dbg !39\n %.extract98 = extractelement <4 x i32> %201, i64 0, !dbg !39\n %.extract100 = extractelement <4 x i32> %201, i64 1, !dbg !39\n %.extract102 = extractelement <4 x i32> %201, i64 2, !dbg !39\n %.extract104 = extractelement <4 x i32> %201, i64 3, !dbg !39\n tail call void asm sideeffect \"@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };\", \"r,r,r,r,l,b\"(i32 %.extract98, i32 %.extract100, i32 %.extract102, i32 %.extract104, ptr addrspace(1) %178, i1 %183) #6, !dbg !39\n ret void, !dbg !65\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)\ndeclare i32 @llvm.smin.i32(i32, i32) #2\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3\n\n; Function Attrs: convergent nocallback nounwind memory(argmem: write)\ndeclare void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) writeonly captures(none), i32, i32, i32, i32) #4\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.bar.warp.sync(i32) #3\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"32\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #3 = { convergent nocallback nounwind }\nattributes #4 = { convergent nocallback nounwind memory(argmem: write) }\nattributes #5 = { convergent nocallback nofree nounwind memory(argmem: read) }\nattributes #6 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"matmul_kernel\", linkageName: \"matmul_kernel\", scope: !1, file: !1, line: 292, type: !5, scopeLine: 292, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !9, !9, !9, !9, !9, !9, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"half\", size: 16, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 311, column: 24, scope: !4)\n!13 = !DILocation(line: 43, column: 17, scope: !14, inlinedAt: !16)\n!14 = distinct !DILexicalBlockFile(scope: !4, file: !15, discriminator: 0)\n!15 = !DIFile(filename: \"standard.py\", directory: \"/python/lib/python3.14/site-packages/triton/language\")\n!16 = !DILocation(line: 312, column: 27, scope: !17)\n!17 = distinct !DILexicalBlockFile(scope: !4, file: !1, discriminator: 0)\n!18 = !DILocation(line: 43, column: 30, scope: !14, inlinedAt: !16)\n!19 = !DILocation(line: 43, column: 17, scope: !14, inlinedAt: !20)\n!20 = !DILocation(line: 313, column: 27, scope: !17)\n!21 = !DILocation(line: 43, column: 30, scope: !14, inlinedAt: !20)\n!22 = !DILocation(line: 315, column: 22, scope: !4)\n!23 = !DILocation(line: 317, column: 33, scope: !4)\n!24 = !DILocation(line: 317, column: 46, scope: !4)\n!25 = !DILocation(line: 318, column: 33, scope: !4)\n!26 = !DILocation(line: 318, column: 27, scope: !4)\n!27 = !DILocation(line: 319, column: 19, scope: !4)\n!28 = !DILocation(line: 319, column: 40, scope: !4)\n!29 = !DILocation(line: 321, column: 23, scope: !4)\n!30 = !DILocation(line: 321, column: 51, scope: !4)\n!31 = !DILocation(line: 321, column: 38, scope: !4)\n!32 = !DILocation(line: 322, column: 23, scope: !4)\n!33 = !DILocation(line: 322, column: 51, scope: !4)\n!34 = !DILocation(line: 322, column: 38, scope: !4)\n!35 = !DILocation(line: 43, column: 17, scope: !14, inlinedAt: !36)\n!36 = !DILocation(line: 328, column: 33, scope: !17)\n!37 = !DILocation(line: 43, column: 30, scope: !14, inlinedAt: !36)\n!38 = !DILocation(line: 328, column: 22, scope: !4)\n!39 = !DILocation(line: 344, column: 21, scope: !4)\n!40 = !DILocation(line: 337, column: 33, scope: !4)\n!41 = !DILocation(line: 322, column: 68, scope: !4)\n!42 = !DILocation(line: 325, column: 36, scope: !4)\n!43 = !DILocation(line: 325, column: 48, scope: !4)\n!44 = !DILocation(line: 325, column: 18, scope: !4)\n!45 = !DILocation(line: 321, column: 68, scope: !4)\n!46 = !DILocation(line: 324, column: 37, scope: !4)\n!47 = !DILocation(line: 324, column: 49, scope: !4)\n!48 = !DILocation(line: 324, column: 18, scope: !4)\n!49 = !DILocation(line: 330, column: 51, scope: !4)\n!50 = !DILocation(line: 330, column: 47, scope: !4)\n!51 = !DILocation(line: 330, column: 43, scope: !4)\n!52 = !DILocation(line: 330, column: 12, scope: !4)\n!53 = !DILocation(line: 333, column: 43, scope: !4)\n!54 = !DILocation(line: 333, column: 12, scope: !4)\n!55 = !DILocation(line: 335, column: 39, scope: !4)\n!56 = !DILocation(line: 336, column: 18, scope: !4)\n!57 = !DILocation(line: 337, column: 18, scope: !4)\n!58 = !DILocation(line: 338, column: 29, scope: !4)\n!59 = !DILocation(line: 342, column: 29, scope: !4)\n!60 = !DILocation(line: 342, column: 17, scope: !4)\n!61 = !DILocation(line: 342, column: 48, scope: !4)\n!62 = !DILocation(line: 343, column: 33, scope: !4)\n!63 = !DILocation(line: 343, column: 58, scope: !4)\n!64 = !DILocation(line: 343, column: 39, scope: !4)\n!65 = !DILocation(line: 344, column: 4, scope: !4)\n","matmul_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.7\n.target sm_90a\n.address_size 64\n\n\t// .globl\tmatmul_kernel // -- Begin function matmul_kernel\n.extern .shared .align 16 .b8 global_smem[];\n // @matmul_kernel\n.visible .entry matmul_kernel(\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_2,\n\t.param .u32 matmul_kernel_param_3,\n\t.param .u32 matmul_kernel_param_4,\n\t.param .u32 matmul_kernel_param_5,\n\t.param .u32 matmul_kernel_param_6,\n\t.param .u32 matmul_kernel_param_7,\n\t.param .u32 matmul_kernel_param_8,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_9,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_10\n)\n.reqntid 32\n{\n\t.reg .pred \t%p<12>;\n\t.reg .b32 \t%r<144>;\n\t.reg .b64 \t%rd<13>;\n\t.loc\t1 292 0 // generate_examples.py:292:0\n$L__func_begin0:\n\t.loc\t1 292 0 // generate_examples.py:292:0\n\n// %bb.0:\n\tld.param.b32 \t%r15, [matmul_kernel_param_8];\n\tld.param.b32 \t%r115, [matmul_kernel_param_5];\n\tld.param.b32 \t%r12, [matmul_kernel_param_4];\n\tld.param.b32 \t%r11, [matmul_kernel_param_3];\n\tld.param.b64 \t%rd4, [matmul_kernel_param_2];\n$L__tmp0:\n\t.loc\t1 311 24 // generate_examples.py:311:24\n\tmov.u32 \t%r16, %ctaid.x;\n$L__tmp1:\n\t.loc\t2 43 17 // standard.py:43:17 @[ generate_examples.py:312:27 ]\n\tadd.s32 \t%r17, %r11, 31;\n\t.loc\t2 43 30 // standard.py:43:30 @[ generate_examples.py:312:27 ]\n\tshr.s32 \t%r18, %r17, 31;\n\tshr.u32 \t%r19, %r18, 27;\n\tadd.s32 \t%r20, %r17, %r19;\n\tshr.s32 \t%r21, %r20, 5;\n$L__tmp2:\n\t.loc\t2 43 17 // standard.py:43:17 @[ generate_examples.py:313:27 ]\n\tadd.s32 \t%r22, %r12, 15;\n\t.loc\t2 43 30 // standard.py:43:30 @[ generate_examples.py:313:27 ]\n\tshr.s32 \t%r23, %r22, 31;\n\tshr.u32 \t%r24, %r23, 28;\n\tadd.s32 \t%r25, %r22, %r24;\n\tshr.s32 \t%r26, %r25, 4;\n$L__tmp3:\n\t.loc\t1 315 22 // generate_examples.py:315:22\n\tdiv.s32 \t%r27, %r16, %r26;\n\t.loc\t1 317 33 // generate_examples.py:317:33\n\tsub.s32 \t%r28, %r21, %r27;\n\t.loc\t1 317 46 // generate_examples.py:317:46\n\tmin.s32 \t%r29, %r28, 1;\n\t.loc\t1 318 33 // generate_examples.py:318:33\n\trem.s32 \t%r30, %r16, %r29;\n\t.loc\t1 318 27 // generate_examples.py:318:27\n\tadd.s32 \t%r31, %r30, %r27;\n\t.loc\t1 319 19 // generate_examples.py:319:19\n\tmul.lo.s32 \t%r32, %r27, %r26;\n\tsub.s32 \t%r33, %r16, %r32;\n\t.loc\t1 319 40 // generate_examples.py:319:40\n\tdiv.s32 \t%r34, %r33, %r29;\n\t.loc\t1 321 23 // generate_examples.py:321:23\n\tshl.b32 \t%r35, %r31, 5;\n\t.loc\t1 321 51 // generate_examples.py:321:51\n\tmov.u32 \t%r1, %tid.x;\n\tbfe.u32 \t%r2, %r1, 1, 4;\n\t.loc\t1 321 38 // generate_examples.py:321:38\n\tor.b32 \t%r3, %r35, %r2;\n\tor.b32 \t%r4, %r3, 16;\n\t.loc\t1 322 23 // generate_examples.py:322:23\n\tshl.b32 \t%r36, %r34, 4;\n\t.loc\t1 322 51 // generate_examples.py:322:51\n\tand.b32 \t%r5, %r1, 1;\n\tshl.b32 \t%r6, %r5, 3;\n\t.loc\t1 322 38 // generate_examples.py:322:38\n\tor.b32 \t%r7, %r36, %r6;\n$L__tmp4:\n\t.loc\t2 43 17 // standard.py:43:17 @[ generate_examples.py:328:33 ]\n\tadd.s32 \t%r37, %r115, 15;\n$L__tmp5:\n\t.loc\t1 328 22 // generate_examples.py:328:22\n\tsetp.gt.s32 \t%p1, %r37, 15;\n\tmov.b32 \t%r114, global_smem;\n\t@%p1 bra \t$L__BB0_2;\n\tbra.uni \t$L__BB0_1;\n$L__BB0_2: // %.lr.ph\n\t.loc\t1 0 22 // generate_examples.py:0:22\n\tld.param.b32 \t%r14, [matmul_kernel_param_7];\n\tld.param.b32 \t%r13, [matmul_kernel_param_6];\n\tld.param.b64 \t%rd3, [matmul_kernel_param_1];\n\tld.param.b64 \t%rd2, [matmul_kernel_param_0];\n\tshr.s32 \t%r38, %r37, 31;\n\tshr.u32 \t%r39, %r38, 28;\n\tadd.s32 \t%r40, %r37, %r39;\n\tshr.s32 \t%r116, %r40, 4;\n\t.loc\t1 337 33 // generate_examples.py:337:33\n\tshl.b32 \t%r41, %r14, 4;\n\t.loc\t1 322 68 // generate_examples.py:322:68\n\trem.s32 \t%r42, %r7, %r12;\n\t.loc\t1 325 48 // generate_examples.py:325:48\n\tmad.lo.s32 \t%r43, %r14, %r2, %r42;\n\t.loc\t1 325 18 // generate_examples.py:325:18\n\tmad.wide.s32 \t%rd10, %r43, 2, %rd3;\n\t.loc\t1 321 68 // generate_examples.py:321:68\n\trem.s32 \t%r44, %r4, %r11;\n\t.loc\t1 324 49 // generate_examples.py:324:49\n\tmad.lo.s32 \t%r45, %r44, %r13, %r6;\n\t.loc\t1 324 18 // generate_examples.py:324:18\n\tmad.wide.s32 \t%rd11, %r45, 2, %rd2;\n\t.loc\t1 321 68 // generate_examples.py:321:68\n\trem.s32 \t%r46, %r3, %r11;\n\t.loc\t1 324 49 // generate_examples.py:324:49\n\tmad.lo.s32 \t%r47, %r46, %r13, %r6;\n\t.loc\t1 324 18 // generate_examples.py:324:18\n\tmad.wide.s32 \t%rd12, %r47, 2, %rd2;\n\tshl.b32 \t%r135, %r1, 4;\n\tand.b32 \t%r48, %r135, 368;\n\tand.b32 \t%r133, %r1, 8;\n\tbfe.s32 \t%r49, %r1, 3, 1;\n\tand.b32 \t%r50, %r49, 144;\n\txor.b32 \t%r51, %r50, %r48;\n\tadd.s32 \t%r71, %r114, %r51;\n\tadd.s32 \t%r72, %r71, 512;\n\tadd.s32 \t%r77, %r71, 1024;\n\tshl.b32 \t%r52, %r1, 5;\n\tand.b32 \t%r53, %r52, 352;\n\tand.b32 \t%r134, %r1, 4;\n\tbfe.s32 \t%r54, %r1, 2, 1;\n\tand.b32 \t%r55, %r54, 144;\n\tand.b32 \t%r56, %r1, 16;\n\tor.b32 \t%r57, %r55, %r53;\n\txor.b32 \t%r58, %r57, %r56;\n\tadd.s32 \t%r8, %r114, %r58;\n\tadd.s32 \t%r9, %r8, 512;\n\txor.b32 \t%r59, %r55, %r56;\n\tadd.s32 \t%r60, %r114, %r53;\n\tadd.s32 \t%r61, %r60, %r59;\n\tadd.s32 \t%r10, %r61, 1024;\n\t.loc\t1 328 22 // generate_examples.py:328:22\n\tmul.wide.s32 \t%rd1, %r41, 2;\n\tmov.b32 \t%r117, 0f00000000;\n\tmov.b32 \t%r118, %r117;\n\tmov.b32 \t%r119, %r117;\n\tmov.b32 \t%r120, %r117;\n\tmov.b32 \t%r121, %r117;\n\tmov.b32 \t%r122, %r117;\n\tmov.b32 \t%r123, %r117;\n\tmov.b32 \t%r124, %r117;\n\tmov.b32 \t%r125, %r117;\n\tmov.b32 \t%r126, %r117;\n\tmov.b32 \t%r127, %r117;\n\tmov.b32 \t%r128, %r117;\n\tmov.b32 \t%r129, %r117;\n\tmov.b32 \t%r130, %r117;\n\tmov.b32 \t%r131, %r117;\n\tmov.b32 \t%r132, %r117;\n$L__BB0_3: // =>This Inner Loop Header: Depth=1\n\t.loc\t1 330 43 // generate_examples.py:330:43\n\tsetp.lt.s32 \t%p2, %r6, %r115;\n\tmov.b32 \t%r66, 0;\n\t.loc\t1 330 12 // generate_examples.py:330:12\n\t// begin inline asm\n\tmov.u32 %r62, %r66;\n\tmov.u32 %r63, %r66;\n\tmov.u32 %r64, %r66;\n\tmov.u32 %r65, %r66;\n\t@%p2 ld.global.v4.b32 { %r62, %r63, %r64, %r65 }, [ %rd12 + 0 ];\n\t// end inline asm\n\t// begin inline asm\n\tmov.u32 %r67, %r66;\n\tmov.u32 %r68, %r66;\n\tmov.u32 %r69, %r66;\n\tmov.u32 %r70, %r66;\n\t@%p2 ld.global.v4.b32 { %r67, %r68, %r69, %r70 }, [ %rd11 + 0 ];\n\t// end inline asm\n\tbar.sync \t0;\n\t// begin inline asm\n\tst.shared.v4.b32 [ %r71 + 0 ], { %r62, %r63, %r64, %r65 };\n\t// end inline asm\n\t// begin inline asm\n\tst.shared.v4.b32 [ %r72 + 0 ], { %r67, %r68, %r69, %r70 };\n\t// end inline asm\n\t.loc\t1 333 43 // generate_examples.py:333:43\n\tsetp.lt.s32 \t%p3, %r2, %r115;\n\t.loc\t1 333 12 // generate_examples.py:333:12\n\t// begin inline asm\n\tmov.u32 %r73, %r66;\n\tmov.u32 %r74, %r66;\n\tmov.u32 %r75, %r66;\n\tmov.u32 %r76, %r66;\n\t@%p3 ld.global.v4.b32 { %r73, %r74, %r75, %r76 }, [ %rd10 + 0 ];\n\t// end inline asm\n\t// begin inline asm\n\tst.shared.v4.b32 [ %r77 + 0 ], { %r73, %r74, %r75, %r76 };\n\t// end inline asm\n\t.loc\t1 330 12 // generate_examples.py:330:12\n\tbar.sync \t0;\n\tldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r78, %r79, %r80, %r81}, [%r8];\n\tldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r86, %r87, %r88, %r89}, [%r9];\n\t.loc\t1 333 12 // generate_examples.py:333:12\n\tldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r82, %r83, %r84, %r85}, [%r10];\n\t.loc\t1 335 39 // generate_examples.py:335:39\n\t// begin inline asm\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r117, %r118, %r119, %r120 }, { %r78, %r79, %r80, %r81 }, { %r82, %r83 }, { %r117, %r118, %r119, %r120 };\n\t// end inline asm\n\t// begin inline asm\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r121, %r122, %r123, %r124 }, { %r78, %r79, %r80, %r81 }, { %r84, %r85 }, { %r121, %r122, %r123, %r124 };\n\t// end inline asm\n\t// begin inline asm\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r125, %r126, %r127, %r128 }, { %r86, %r87, %r88, %r89 }, { %r82, %r83 }, { %r125, %r126, %r127, %r128 };\n\t// end inline asm\n\t// begin inline asm\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r129, %r130, %r131, %r132 }, { %r86, %r87, %r88, %r89 }, { %r84, %r85 }, { %r129, %r130, %r131, %r132 };\n\t// end inline asm\n\t.loc\t1 336 18 // generate_examples.py:336:18\n\tadd.s64 \t%rd12, %rd12, 32;\n\tadd.s64 \t%rd11, %rd11, 32;\n\t.loc\t1 328 22 // generate_examples.py:328:22\n\tadd.s64 \t%rd10, %rd10, %rd1;\n\tadd.s32 \t%r116, %r116, -1;\n\tadd.s32 \t%r115, %r115, -16;\n\tsetp.ne.b32 \t%p4, %r116, 0;\n\t@%p4 bra \t$L__BB0_3;\n// %bb.4: // %._crit_edge.loopexit\n\t.loc\t1 338 29 // generate_examples.py:338:29\n\tcvt.rn.f16x2.f32 \t%r143, %r132, %r131;\n\tcvt.rn.f16x2.f32 \t%r142, %r130, %r129;\n\tcvt.rn.f16x2.f32 \t%r141, %r128, %r127;\n\tcvt.rn.f16x2.f32 \t%r140, %r126, %r125;\n\tcvt.rn.f16x2.f32 \t%r139, %r124, %r123;\n\tcvt.rn.f16x2.f32 \t%r138, %r122, %r121;\n\tcvt.rn.f16x2.f32 \t%r137, %r120, %r119;\n\tcvt.rn.f16x2.f32 \t%r136, %r118, %r117;\n\tbra.uni \t$L__BB0_5;\n$L__BB0_1: // %.._crit_edge_crit_edge\n\t.loc\t1 344 21 // generate_examples.py:344:21\n\tshl.b32 \t%r135, %r1, 4;\n\tand.b32 \t%r134, %r1, 4;\n\tand.b32 \t%r133, %r1, 8;\n\tmov.b32 \t%r136, 0;\n\tmov.b32 \t%r137, %r136;\n\tmov.b32 \t%r138, %r136;\n\tmov.b32 \t%r139, %r136;\n\tmov.b32 \t%r140, %r136;\n\tmov.b32 \t%r141, %r136;\n\tmov.b32 \t%r142, %r136;\n\tmov.b32 \t%r143, %r136;\n$L__BB0_5: // %._crit_edge\n\t.loc\t1 342 29 // generate_examples.py:342:29\n\tmul.lo.s32 \t%r98, %r3, %r15;\n\tmul.lo.s32 \t%r99, %r15, %r4;\n\t.loc\t1 342 17 // generate_examples.py:342:17\n\tmad.wide.s32 \t%rd7, %r98, 2, %rd4;\n\tmad.wide.s32 \t%rd8, %r99, 2, %rd4;\n\t.loc\t1 342 48 // generate_examples.py:342:48\n\tmul.wide.s32 \t%rd9, %r7, 2;\n\tadd.s64 \t%rd5, %rd7, %rd9;\n\tadd.s64 \t%rd6, %rd8, %rd9;\n\t.loc\t1 343 33 // generate_examples.py:343:33\n\tsetp.lt.s32 \t%p7, %r3, %r11;\n\tsetp.lt.s32 \t%p8, %r4, %r11;\n\t.loc\t1 343 58 // generate_examples.py:343:58\n\tsetp.lt.s32 \t%p9, %r7, %r12;\n\t.loc\t1 343 39 // generate_examples.py:343:39\n\tand.pred \t%p5, %p7, %p9;\n\tand.pred \t%p6, %p8, %p9;\n\t.loc\t1 344 21 // generate_examples.py:344:21\n\tbar.sync \t0;\n\tand.b32 \t%r100, %r135, 176;\n\tsetp.eq.b32 \t%p10, %r134, 0;\n\tselp.b32 \t%r101, 0, 320, %p10;\n\tshl.b32 \t%r102, %r1, 2;\n\tand.b32 \t%r103, %r102, 64;\n\tor.b32 \t%r104, %r101, %r100;\n\txor.b32 \t%r105, %r104, %r103;\n\tadd.s32 \t%r106, %r114, %r105;\n\tstmatrix.sync.aligned.m8n8.x4.shared.b16 [%r106], {%r136, %r137, %r138, %r139};\n\tbar.warp.sync \t-1;\n\tshl.b32 \t%r107, %r1, 3;\n\tand.b32 \t%r108, %r107, 176;\n\tshl.b32 \t%r109, %r5, 6;\n\tsetp.eq.b32 \t%p11, %r133, 0;\n\tselp.b32 \t%r110, 0, 320, %p11;\n\tor.b32 \t%r111, %r108, %r109;\n\txor.b32 \t%r112, %r111, %r110;\n\tadd.s32 \t%r113, %r114, %r112;\n\tld.shared.v4.b32 \t{%r90, %r91, %r92, %r93}, [%r113];\n\tbar.warp.sync \t-1;\n\tstmatrix.sync.aligned.m8n8.x4.shared.b16 [%r106], {%r140, %r141, %r142, %r143};\n\tbar.warp.sync \t-1;\n\tld.shared.v4.b32 \t{%r94, %r95, %r96, %r97}, [%r113];\n\t// begin inline asm\n\t@%p5 st.global.v4.b32 [ %rd5 + 0 ], { %r90, %r91, %r92, %r93 };\n\t// end inline asm\n\t// begin inline asm\n\t@%p6 st.global.v4.b32 [ %rd6 + 0 ], { %r94, %r95, %r96, %r97 };\n\t// end inline asm\n\t.loc\t1 344 4 // generate_examples.py:344:4\n\tret;\n$L__tmp6:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.file\t2 \"/python/lib/python3.14/site-packages/triton/language/standard.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 1 // DW_CHILDREN_yes\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 2 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 0 // DW_CHILDREN_no\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 32 // DW_AT_inline\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 3 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 1 // DW_CHILDREN_yes\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 4 // Abbreviation Code\n.b8 29 // DW_TAG_inlined_subroutine\n.b8 0 // DW_CHILDREN_no\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 88 // DW_AT_call_file\n.b8 11 // DW_FORM_data1\n.b8 89 // DW_AT_call_line\n.b8 5 // DW_FORM_data2\n.b8 87 // DW_AT_call_column\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 245 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0xee DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 45\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 45\n.b8 114\n.b8 101\n.b8 103\n.b8 101\n.b8 110\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n.b8 2 // Abbrev [2] 0x87:0x10 DW_TAG_subprogram\n.b8 109 // DW_AT_name\n.b8 97\n.b8 116\n.b8 109\n.b8 117\n.b8 108\n.b8 95\n.b8 107\n.b8 101\n.b8 114\n.b8 110\n.b8 101\n.b8 108\n.b8 0\n.b8 1 // DW_AT_inline\n.b8 3 // Abbrev [3] 0x97:0x61 DW_TAG_subprogram\n.b64 $L__func_begin0 // DW_AT_low_pc\n.b64 $L__func_end0 // DW_AT_high_pc\n.b32 135 // DW_AT_abstract_origin\n.b8 4 // Abbrev [4] 0xac:0x19 DW_TAG_inlined_subroutine\n.b32 135 // DW_AT_abstract_origin\n.b64 $L__tmp1 // DW_AT_low_pc\n.b64 $L__tmp2 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 56 // DW_AT_call_line\n.b8 1\n.b8 27 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xc5:0x19 DW_TAG_inlined_subroutine\n.b32 135 // DW_AT_abstract_origin\n.b64 $L__tmp2 // DW_AT_low_pc\n.b64 $L__tmp3 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 57 // DW_AT_call_line\n.b8 1\n.b8 27 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xde:0x19 DW_TAG_inlined_subroutine\n.b32 135 // DW_AT_abstract_origin\n.b64 $L__tmp4 // DW_AT_low_pc\n.b64 $L__tmp5 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 72 // DW_AT_call_line\n.b8 1\n.b8 33 // DW_AT_call_column\n.b8 0 // End Of Children Mark\n.b8 0 // End Of Children Mark\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","matmul_kernel.json":"{\"hash\": \"f7336a625a0c6f924e5fee4b7063c21567b390bb49fc8cc4fd07cd7d06898f5e\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 1, \"num_ctas\": 1, \"num_stages\": 1, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.7.0\", \"tensordesc_meta\": [], \"shared\": 1536, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"matmul_kernel\"}","matmul_kernel.sass":"Function:matmul_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*0010*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 311\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*0020*/ S2R R7, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 96\n /*0030*/ UMOV UR4, 0x400 ;\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0040*/ ULDC.64 UR6, c[0x0][0x208] ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 92\n /*0050*/ LDC R58, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 96\n /*0060*/ S2UR UR5, SR_CgaCtaId ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 53\n /*0070*/ VIADD R0, R9, 0xf ;\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*0080*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 92\n /*0090*/ VIADD R58, R58, 0xf ;\n\t//## File \".nv_debug_ptx_txt\", line 57\n /*00a0*/ LEA.HI R3, R3, R0, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*00b0*/ IABS R10, R7 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 58\n /*00c0*/ SHF.R.S32.HI R4, RZ, 0x4, R3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 96\n /*00d0*/ ULEA UR4, UR5, UR4, 0x18 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*00e0*/ IABS R5, R4.reuse ;\n /*00f0*/ IABS R12, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 39\n /*0100*/ ULDC UR5, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*0110*/ I2F.RP R0, R5 ;\n /*0120*/ MUFU.RCP R0, R0 ;\n /*0130*/ VIADD R2, R0, 0xffffffe ;\n /*0140*/ IMAD.MOV R0, RZ, RZ, -R12 ;\n /*0150*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0160*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0170*/ IMAD.MOV R6, RZ, RZ, -R3 ;\n /*0180*/ IMAD R11, R6, R5, RZ ;\n /*0190*/ IMAD.MOV.U32 R6, RZ, RZ, R10 ;\n /*01a0*/ IMAD.HI.U32 R3, R3, R11, R2 ;\n /*01b0*/ IMAD.HI.U32 R3, R3, R6, RZ ;\n /*01c0*/ IMAD R0, R3, R0, R6 ;\n /*01d0*/ ISETP.GT.U32.AND P1, PT, R5, R0, PT ;\n /*01e0*/ @!P1 IMAD.IADD R0, R0, 0x1, -R5 ;\n /*01f0*/ @!P1 VIADD R3, R3, 0x1 ;\n /*0200*/ ISETP.NE.AND P1, PT, R4, RZ, PT ;\n /*0210*/ ISETP.GE.U32.AND P0, PT, R0, R5, PT ;\n /*0220*/ LOP3.LUT R0, R7, R4, RZ, 0x3c, !PT ;\n /*0230*/ ISETP.GE.AND P2, PT, R0, RZ, PT ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*0240*/ VIADD R0, R8, 0x1f ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*0250*/ @P0 VIADD R3, R3, 0x1 ;\n /*0260*/ IMAD.MOV.U32 R12, RZ, RZ, R3 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 47\n /*0270*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*0280*/ @!P2 IMAD.MOV R12, RZ, RZ, -R12 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43\n\t//## File \".nv_debug_ptx_txt\", line 49\n /*0290*/ LEA.HI R3, R3, R0, RZ, 0x5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 61\n /*02a0*/ @!P1 LOP3.LUT R12, RZ, R4, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*02b0*/ ISETP.GE.AND P2, PT, R7, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*02c0*/ LEA.HI.SX32 R3, R3, -R12, 0x1b ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 72\n /*02d0*/ IMAD.MOV R5, RZ, RZ, -R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*02e0*/ VIMNMX R15, R3, 0x1, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 72\n /*02f0*/ IMAD R4, R4, R5, R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0300*/ IABS R11, R15.reuse ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0310*/ IABS R14, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0320*/ I2F.RP R0, R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0330*/ LOP3.LUT R4, R4, R15, RZ, 0x3c, !PT ;\n /*0340*/ ISETP.GE.AND P4, PT, R4, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0350*/ MUFU.RCP R0, R0 ;\n /*0360*/ VIADD R2, R0, 0xffffffe ;\n /*0370*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0380*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0390*/ IMAD.MOV R10, RZ, RZ, -R3 ;\n /*03a0*/ IMAD R5, R10, R11, RZ ;\n /*03b0*/ IABS R10, R15 ;\n /*03c0*/ IMAD.HI.U32 R3, R3, R5, R2 ;\n /*03d0*/ IMAD.MOV R13, RZ, RZ, -R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*03e0*/ IMAD.MOV.U32 R10, RZ, RZ, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*03f0*/ IMAD.HI.U32 R0, R3, R6, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0400*/ IMAD.HI.U32 R5, R3, R10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0410*/ IMAD R2, R0, R13.reuse, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0420*/ IMAD R3, R5, R13, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*0430*/ S2R R0, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0440*/ ISETP.GT.U32.AND P0, PT, R11.reuse, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0450*/ ISETP.GT.U32.AND P3, PT, R11, R3, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0460*/ @!P0 IMAD.IADD R2, R2, 0x1, -R11.reuse ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0470*/ @!P3 IMAD.IADD R3, R3, 0x1, -R11 ;\n /*0480*/ @!P3 VIADD R5, R5, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0490*/ ISETP.GT.U32.AND P1, PT, R11, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*04a0*/ ISETP.GE.U32.AND P0, PT, R3, R11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*04b0*/ LOP3.LUT R3, RZ, R15, RZ, 0x33, !PT ;\n /*04c0*/ @!P1 IMAD.IADD R2, R2, 0x1, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*04d0*/ @P0 VIADD R5, R5, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*04e0*/ ISETP.NE.AND P0, PT, R15, RZ, PT ;\n /*04f0*/ @!P2 IMAD.MOV R2, RZ, RZ, -R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0500*/ @!P4 IMAD.MOV R5, RZ, RZ, -R5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*0510*/ SEL R2, R3.reuse, R2, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0520*/ SEL R7, R3, R5, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 95\n /*0530*/ ISETP.GT.AND P0, PT, R58, 0xf, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0540*/ IMAD.IADD R6, R12, 0x1, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0550*/ SHF.R.U32.HI R3, RZ, 0x1, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 86\n /*0560*/ LOP3.LUT R2, R0, 0x1, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0570*/ IMAD.SHL.U32 R6, R6, 0x20, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0580*/ SGXT.U32 R3, R3, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 87\n /*0590*/ IMAD.SHL.U32 R4, R2, 0x8, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*05a0*/ LOP3.LUT R5, R6, R3, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 89\n /*05b0*/ IMAD R6, R7, 0x10, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 82\n /*05c0*/ LOP3.LUT R7, R5, 0x10, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 97\n /*05d0*/ @P0 BRA `(.L_x_0) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 250\n /*05e0*/ IMAD.SHL.U32 R48, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 251\n /*05f0*/ LOP3.LUT R49, R0.reuse, 0x4, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 260\n /*0600*/ CS2R R12, SRZ ;\n\t//## File \".nv_debug_ptx_txt\", line 252\n /*0610*/ LOP3.LUT R52, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 260\n /*0620*/ CS2R R14, SRZ ;\n /*0630*/ CS2R R20, SRZ ;\n /*0640*/ CS2R R22, SRZ ;\n /*0650*/ BRA `(.L_x_1) ;\n.L_x_0:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0660*/ IABS R17, R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0670*/ LDC R60, c[0x0][0x238] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0680*/ IABS R16, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0690*/ ULDC UR10, c[0x0][0x234] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*06a0*/ I2F.RP R14, R17 ;\n /*06b0*/ ISETP.NE.AND P3, PT, R9, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*06c0*/ ULDC.64 UR8, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 132\n /*06d0*/ SHF.R.S32.HI R23, RZ, 0x3, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*06e0*/ ULDC.64 UR12, c[0x0][0x210] ;\n\t//## File \".nv_debug_ptx_txt\", line 129\n /*06f0*/ IMAD.SHL.U32 R48, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 131\n /*0700*/ LOP3.LUT R52, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0710*/ I2F.RP R15, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 132\n /*0720*/ SGXT R23, R23, 0x1 ;\n\t//## File \".nv_debug_ptx_txt\", line 140\n /*0730*/ LOP3.LUT R49, R0, 0x4, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 133\n /*0740*/ LOP3.LUT R53, R23, 0x90, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0750*/ CS2R R22, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0760*/ MUFU.RCP R14, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 134\n /*0770*/ LOP3.LUT R53, R53, 0x170, R48, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0780*/ MUFU.RCP R15, R15 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0790*/ VIADD R10, R14, 0xffffffe ;\n /*07a0*/ IABS R14, R6 ;\n /*07b0*/ F2I.FTZ.U32.TRUNC.NTZ R11, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*07c0*/ VIADD R12, R15, 0xffffffe ;\n /*07d0*/ F2I.FTZ.U32.TRUNC.NTZ R13, R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*07e0*/ IMAD.MOV.U32 R10, RZ, RZ, RZ ;\n /*07f0*/ IMAD.MOV R18, RZ, RZ, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0800*/ IMAD.MOV.U32 R12, RZ, RZ, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0810*/ IMAD R19, R18, R17, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0820*/ IMAD.MOV R21, RZ, RZ, -R13 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0830*/ IMAD.HI.U32 R11, R11, R19, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0840*/ IABS R10, R7 ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0850*/ IABS R19, R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0860*/ IMAD R15, R21, R16, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0870*/ CS2R R20, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0880*/ IMAD.HI.U32 R11, R11, R14, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0890*/ IMAD.HI.U32 R12, R13, R15, R12 ;\n /*08a0*/ IMAD.MOV.U32 R13, RZ, RZ, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*08b0*/ IMAD.MOV R11, RZ, RZ, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*08c0*/ IMAD.HI.U32 R10, R12, R13, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*08d0*/ IMAD.HI.U32 R12, R12, R19, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*08e0*/ IMAD.MOV R15, RZ, RZ, -R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*08f0*/ IMAD.MOV R18, RZ, RZ, -R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0900*/ IMAD R12, R16.reuse, R15, R13 ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0910*/ IMAD R13, R16.reuse, R18, R19 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0920*/ CS2R R18, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0930*/ IMAD R10, R17, R11, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0940*/ ISETP.GT.U32.AND P1, PT, R16.reuse, R12, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0950*/ ISETP.GT.U32.AND P2, PT, R16, R13, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0960*/ ISETP.GT.U32.AND P0, PT, R17, R10, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 105\n /*0970*/ SHF.R.S32.HI R11, RZ, 0x1f, R58 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 141\n /*0980*/ SHF.R.S32.HI R14, RZ, 0x2, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 107\n /*0990*/ LEA.HI R58, R11, R58, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*09a0*/ LOP3.LUT R11, RZ, R8, RZ, 0x33, !PT ;\n /*09b0*/ @!P1 IMAD.IADD R12, R12, 0x1, -R16.reuse ;\n /*09c0*/ ISETP.GE.AND P1, PT, R7, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*09d0*/ @!P2 IMAD.IADD R13, R13, 0x1, -R16 ;\n /*09e0*/ ISETP.GE.AND P2, PT, R5, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*09f0*/ @!P0 IMAD.IADD R10, R10, 0x1, -R17 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0a00*/ ISETP.GT.U32.AND P5, PT, R16, R12, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0a10*/ ISETP.GT.U32.AND P6, PT, R16, R13, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0a20*/ ISETP.GT.U32.AND P4, PT, R17, R10, PT ;\n /*0a30*/ ISETP.GE.AND P0, PT, R6, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 108\n /*0a40*/ SHF.R.S32.HI R58, RZ, 0x4, R58 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0a50*/ @!P5 IMAD.IADD R12, R12, 0x1, -R16.reuse ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0a60*/ @!P6 IMAD.IADD R13, R13, 0x1, -R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0a70*/ @!P4 IMAD.IADD R10, R10, 0x1, -R17 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0a80*/ ISETP.NE.AND P4, PT, R8, RZ, PT ;\n /*0a90*/ @!P1 IMAD.MOV R12, RZ, RZ, -R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0aa0*/ CS2R R16, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0ab0*/ @!P2 IMAD.MOV R13, RZ, RZ, -R13 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0ac0*/ @!P0 IMAD.MOV R10, RZ, RZ, -R10 ;\n /*0ad0*/ @!P3 LOP3.LUT R10, RZ, R9, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 138\n /*0ae0*/ IMAD.SHL.U32 R8, R0, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0af0*/ SEL R9, R11, R12, !P4 ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0b00*/ SEL R11, R11, R13, !P4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0b10*/ IMAD R10, R3, R60, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 141\n /*0b20*/ SGXT R12, R14, 0x1 ;\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0b30*/ IMAD R9, R9, UR10, R4.reuse ;\n\t//## File \".nv_debug_ptx_txt\", line 139\n /*0b40*/ LOP3.LUT R8, R8, 0x160, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0b50*/ IMAD R11, R11, UR10, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 142\n /*0b60*/ LOP3.LUT R13, R12, 0x90, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0b70*/ CS2R R14, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*0b80*/ LEA R54, P1, R9, UR12, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 337\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0b90*/ IMAD.SHL.U32 R60, R60, 0x10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 148\n /*0ba0*/ LOP3.LUT R55, R13, 0x10, R0, 0x78, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 128\n /*0bb0*/ LEA R56, P2, R11, UR12, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0bc0*/ LEA R38, P0, R10, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 144\n /*0bd0*/ LOP3.LUT R57, R13, R8, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0be0*/ CS2R R12, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 150\n /*0bf0*/ IADD3 R55, R55, UR4, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*0c00*/ LEA.HI.X.SX32 R59, R9, UR13, 0x1, P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0c10*/ CS2R R8, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 128\n /*0c20*/ LEA.HI.X.SX32 R61, R11, UR13, 0x1, P2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0c30*/ LEA.HI.X.SX32 R39, R10, UR9, 0x1, P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 169\n /*0c40*/ CS2R R10, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 145\n /*0c50*/ LOP3.LUT R57, R57, 0x10, R0, 0x78, !PT ;\n.L_x_2:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 172\n /*0c60*/ ISETP.GE.AND P1, PT, R4, UR5, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 180\n /*0c70*/ CS2R R24, SRZ ;\n /*0c80*/ CS2R R26, SRZ ;\n\t//## File \".nv_debug_ptx_txt\", line 187\n /*0c90*/ CS2R R28, SRZ ;\n /*0ca0*/ CS2R R30, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 197\n /*0cb0*/ ISETP.GE.AND P0, PT, R3, UR5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 180\n /*0cc0*/ @!P1 IMAD.MOV.U32 R36, RZ, RZ, R56 ;\n /*0cd0*/ @!P1 IMAD.MOV.U32 R37, RZ, RZ, R61 ;\n\t//## File \".nv_debug_ptx_txt\", line 187\n /*0ce0*/ @!P1 IMAD.MOV.U32 R50, RZ, RZ, R54 ;\n /*0cf0*/ @!P1 IMAD.MOV.U32 R51, RZ, RZ, R59 ;\n\t//## File \".nv_debug_ptx_txt\", line 180\n /*0d00*/ @!P1 LDG.E.128 R24, desc[UR6][R36.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 204\n /*0d10*/ CS2R R32, SRZ ;\n /*0d20*/ CS2R R34, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 187\n /*0d30*/ @!P1 LDG.E.128 R28, desc[UR6][R50.64] ;\n\t//## File \".nv_debug_ptx_txt\", line 189\n /*0d40*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n /*0d50*/ IMAD.MOV.U32 R50, RZ, RZ, R38 ;\n /*0d60*/ IMAD.MOV.U32 R51, RZ, RZ, R39 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 204\n /*0d70*/ @!P0 LDG.E.128 R32, desc[UR6][R50.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 233\n /*0d80*/ VIADD R58, R58, 0xffffffff ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 229\n /*0d90*/ IADD3 R56, P1, R56, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 234\n /*0da0*/ UIADD3 UR5, UR5, -0x10, URZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 230\n /*0db0*/ IADD3 R54, P2, R54, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0dc0*/ ISETP.NE.U32.AND P0, PT, R58, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 229\n /*0dd0*/ IMAD.X R61, RZ, RZ, R61, P1 ;\n\t//## File \".nv_debug_ptx_txt\", line 230\n /*0de0*/ IMAD.X R59, RZ, RZ, R59, P2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 191\n /*0df0*/ STS.128 [R53+UR4], R24 ;\n\t//## File \".nv_debug_ptx_txt\", line 194\n /*0e00*/ STS.128 [R53+UR4+0x200], R28 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 207\n /*0e10*/ STS.128 [R53+UR4+0x400], R32 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 210\n /*0e20*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 211\n /*0e30*/ LDSM.16.M88.4 R36, [R57+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 214\n /*0e40*/ LDSM.16.MT88.4 R40, [R55+0x400] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 212\n /*0e50*/ LDSM.16.M88.4 R44, [R57+UR4+0x200] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 217\n /*0e60*/ HMMA.16816.F32 R8, R36.reuse, R40, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 220\n /*0e70*/ HMMA.16816.F32 R12, R36, R42, R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 223\n /*0e80*/ HMMA.16816.F32 R16, R44, R40, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 232\n /*0e90*/ IMAD.WIDE R38, R60, 0x2, R50 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 226\n /*0ea0*/ HMMA.16816.F32 R20, R44, R42, R20 ;\n /*0eb0*/ NOP ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 236\n /*0ec0*/ @P0 BRA `(.L_x_2) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 338\n\t//## File \".nv_debug_ptx_txt\", line 239\n /*0ed0*/ F2FP.F16.F32.PACK_AB R23, R23, R22 ;\n\t//## File \".nv_debug_ptx_txt\", line 243\n /*0ee0*/ F2FP.F16.F32.PACK_AB R15, R15, R14 ;\n\t//## File \".nv_debug_ptx_txt\", line 240\n /*0ef0*/ F2FP.F16.F32.PACK_AB R22, R21, R20 ;\n\t//## File \".nv_debug_ptx_txt\", line 244\n /*0f00*/ F2FP.F16.F32.PACK_AB R14, R13, R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 241\n /*0f10*/ F2FP.F16.F32.PACK_AB R21, R19, R18 ;\n\t//## File \".nv_debug_ptx_txt\", line 242\n /*0f20*/ F2FP.F16.F32.PACK_AB R20, R17, R16 ;\n\t//## File \".nv_debug_ptx_txt\", line 245\n /*0f30*/ F2FP.F16.F32.PACK_AB R13, R11, R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 246\n /*0f40*/ F2FP.F16.F32.PACK_AB R12, R9, R8 ;\n.L_x_1:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 281\n /*0f50*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 283\n /*0f60*/ ISETP.NE.U32.AND P0, PT, R49, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 285\n /*0f70*/ IMAD.SHL.U32 R4, R0.reuse, 0x4, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 292\n /*0f80*/ IMAD.SHL.U32 R0, R0, 0x8, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 284\n /*0f90*/ SEL R3, RZ, 0x140, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 263\n /*0fa0*/ ULDC UR5, c[0x0][0x23c] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 295\n /*0fb0*/ ISETP.NE.U32.AND P0, PT, R52, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 276\n /*0fc0*/ ULDC.64 UR10, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 287\n /*0fd0*/ LOP3.LUT R3, R3, 0xb0, R48, 0xf8, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 266\n /*0fe0*/ ULDC.64 UR8, c[0x0][0x220] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 293\n /*0ff0*/ LOP3.LUT R9, R0, 0xb0, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 264\n /*1000*/ IMAD R8, R7, UR5, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 288\n /*1010*/ LOP3.LUT R3, R3, 0x40, R4, 0x78, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 296\n /*1020*/ SEL R0, RZ, 0x140, !P0 ;\n\t//## File \".nv_debug_ptx_txt\", line 297\n /*1030*/ IMAD R11, R2, 0x40, R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 276\n /*1040*/ ISETP.GE.AND P0, PT, R6, UR11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 289\n /*1050*/ VIADD R9, R3, UR4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 267\n /*1060*/ LEA R4, P3, R8, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 298\n /*1070*/ LOP3.LUT R11, R11, R0, RZ, 0x3c, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 263\n /*1080*/ IMAD R0, R5.reuse, UR5, RZ ;\n /*1090*/ ISETP.LT.AND P1, PT, R5, UR10, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 290\n /*10a0*/ STSM.16.M88.4 [R9], R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 291\n /*10b0*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 300\n /*10c0*/ LDS.128 R16, [R11+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 266\n /*10d0*/ LEA R2, P2, R0.reuse, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 301\n /*10e0*/ NOP ;\n /*10f0*/ ISETP.LT.AND P0, PT, R7, UR10, !P0 ;\n\t//## File \".nv_debug_ptx_txt\", line 302\n /*1100*/ STSM.16.M88.4 [R9], R20 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 266\n /*1110*/ LEA.HI.X.SX32 R3, R0, UR9, 0x1, P2 ;\n\t//## File \".nv_debug_ptx_txt\", line 267\n /*1120*/ LEA.HI.X.SX32 R5, R8, UR9, 0x1, P3 ;\n\t//## File \".nv_debug_ptx_txt\", line 270\n /*1130*/ IMAD.WIDE R2, R6, 0x2, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 303\n /*1140*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 306\n /*1150*/ @P1 STG.E.128 desc[UR6][R2.64], R16 ;\n\t//## File \".nv_debug_ptx_txt\", line 309\n /*1160*/ @!P0 EXIT ;\n\t//## File \".nv_debug_ptx_txt\", line 304\n /*1170*/ LDS.128 R8, [R11+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 271\n /*1180*/ IMAD.WIDE R4, R6, 0x2, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 309\n /*1190*/ STG.E.128 desc[UR6][R4.64], R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 312\n /*11a0*/ EXIT ;\n.L_x_3:\n /*11b0*/ BRA `(.L_x_3);\n /*11c0*/ NOP;\n /*11d0*/ NOP;\n /*11e0*/ NOP;\n /*11f0*/ NOP;\n /*1200*/ NOP;\n /*1210*/ NOP;\n /*1220*/ NOP;\n /*1230*/ NOP;\n /*1240*/ NOP;\n /*1250*/ NOP;\n /*1260*/ NOP;\n /*1270*/ NOP;\n.L_x_4:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":266,"end_line":344,"code":" @triton.autotune(\n configs=[\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 16,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 32,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n ],\n key=[\"M\", \"N\", \"K\"],\n )\n @triton.jit\n def matmul_kernel(\n a,\n b,\n c,\n M,\n N,\n K,\n stride_am,\n stride_ak,\n stride_bk,\n stride_bn,\n stride_cm,\n stride_cn,\n BLOCK_SIZE_M: tl.constexpr,\n BLOCK_SIZE_N: tl.constexpr,\n BLOCK_SIZE_K: tl.constexpr,\n GROUP_SIZE_M: tl.constexpr,\n ACTIVATION: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)\n num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)\n num_pid_in_group = GROUP_SIZE_M * num_pid_n\n group_id = pid // num_pid_in_group\n first_pid_m = group_id * GROUP_SIZE_M\n group_size = min(num_pid_m - first_pid_m, GROUP_SIZE_M)\n pid_m = first_pid_m + (pid % group_size)\n pid_n = (pid % num_pid_in_group) // group_size\n\n offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M\n offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N\n offs_k = tl.arange(0, BLOCK_SIZE_K)\n a_ptrs = a + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)\n b_ptrs = b + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)\n\n accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)\n for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):\n a_block = tl.load(\n a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0\n )\n b_block = tl.load(\n b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0\n )\n accumulator += tl.dot(a_block, b_block)\n a_ptrs += BLOCK_SIZE_K * stride_ak\n b_ptrs += BLOCK_SIZE_K * stride_bk\n c_block = accumulator.to(tl.float16)\n\n offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)\n offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)\n c_ptrs = c + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]\n c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)\n tl.store(c_ptrs, c_block, mask=c_mask)\n"}}} -{"event_type":"compilation","pid":1724463,"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":429,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op(x, y, z, scale_factor=1.0, activation=\"none\").sum()"},{"line":404,"name":"fused_op","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":739,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":884,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:23.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"1aa7bd5d4058894489843d4e8c2bfc6c480457536a3cc8e94bc4553e36fb837e","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"fused_op_kernel","num_ctas":1,"num_stages":3,"num_warps":4,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":0,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[],"(6,)":[],"(7,)":[]},"src_cache_key":"f2e5ae6c41d3bfc26af60eada8fff38a0b318cffe77df20b32d117346110ff07","src_constants":{"(6,)":"none","(7,)":8},"times":{"ir_initialization":1602,"lowering_stages":[],"store_results":0}},"file_path":{"fused_op_kernel.source":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.source","fused_op_kernel.ttir":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.ttir","fused_op_kernel.ttgir":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.ttgir","fused_op_kernel.llir":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.llir","fused_op_kernel.ptx":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.ptx","fused_op_kernel.cubin":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.cubin","fused_op_kernel.json":"/home/user/.triton/cache/DKT32XKALCEUJCMEHVHIYK74NREAIV2TNI6MR2KLYRKT4NX3QN7A/fused_op_kernel.json"},"file_content":{"fused_op_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:0)\n#loc19 = loc(\"a_ptr\"(#loc))\n#loc20 = loc(\"b_ptr\"(#loc))\n#loc21 = loc(\"c_ptr\"(#loc))\n#loc22 = loc(\"output_ptr\"(#loc))\n#loc23 = loc(\"n_elements\"(#loc))\n#loc24 = loc(\"scale_factor\"(#loc))\nmodule {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc25)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc26)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc27)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32> loc(#loc28)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32> loc(#loc28)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32> loc(#loc29)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32> loc(#loc29)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc30)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc30)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr> loc(#loc31)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc32)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc32)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr> loc(#loc33)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc34)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc34)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr> loc(#loc35)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32> loc(#loc36)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32> loc(#loc37)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32> loc(#loc37)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32> loc(#loc38)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc16)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc16)\n tt.store %1, %result_12, %mask_3 : tensor<8x!tt.ptr> loc(#loc17)\n tt.return loc(#loc18)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:24)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:20)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:46)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:33)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:24)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:16)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:24)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:16)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:24)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:16)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:17)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:21)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:36)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:26)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:35)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:4)\n#loc25 = loc(\"pid\"(#loc2))\n#loc26 = loc(\"offsets\"(#loc3))\n#loc27 = loc(\"offsets\"(#loc4))\n#loc28 = loc(\"offsets\"(#loc5))\n#loc29 = loc(\"mask\"(#loc6))\n#loc30 = loc(\"a\"(#loc7))\n#loc31 = loc(\"a\"(#loc8))\n#loc32 = loc(\"b\"(#loc9))\n#loc33 = loc(\"b\"(#loc10))\n#loc34 = loc(\"c\"(#loc11))\n#loc35 = loc(\"c\"(#loc12))\n#loc36 = loc(\"result\"(#loc13))\n#loc37 = loc(\"result\"(#loc14))\n#loc38 = loc(\"result\"(#loc15))\n","fused_op_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:0)\n#loc19 = loc(\"a_ptr\"(#loc))\n#loc20 = loc(\"b_ptr\"(#loc))\n#loc21 = loc(\"c_ptr\"(#loc))\n#loc22 = loc(\"output_ptr\"(#loc))\n#loc23 = loc(\"n_elements\"(#loc))\n#loc24 = loc(\"scale_factor\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 4 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc25)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc26)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #blocked> loc(#loc27)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32, #blocked> loc(#loc28)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32, #blocked> loc(#loc28)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32, #blocked> loc(#loc29)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32, #blocked> loc(#loc29)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc30)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc30)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc31)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc32)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc32)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc33)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc34)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc34)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc35)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32, #blocked> loc(#loc36)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32, #blocked> loc(#loc37)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32, #blocked> loc(#loc37)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32, #blocked> loc(#loc38)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc16)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc16)\n tt.store %1, %result_12, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc17)\n tt.return loc(#loc18)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:24)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:20)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:46)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:33)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:24)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:16)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:24)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:16)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:24)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:16)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:17)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:21)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:36)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:26)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:35)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:4)\n#loc25 = loc(\"pid\"(#loc2))\n#loc26 = loc(\"offsets\"(#loc3))\n#loc27 = loc(\"offsets\"(#loc4))\n#loc28 = loc(\"offsets\"(#loc5))\n#loc29 = loc(\"mask\"(#loc6))\n#loc30 = loc(\"a\"(#loc7))\n#loc31 = loc(\"a\"(#loc8))\n#loc32 = loc(\"b\"(#loc9))\n#loc33 = loc(\"b\"(#loc10))\n#loc34 = loc(\"c\"(#loc11))\n#loc35 = loc(\"c\"(#loc12))\n#loc36 = loc(\"result\"(#loc13))\n#loc37 = loc(\"result\"(#loc14))\n#loc38 = loc(\"result\"(#loc15))\n","fused_op_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @fused_op_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, i32 %4, float %5, ptr addrspace(1) readnone captures(none) %6, ptr addrspace(1) readnone captures(none) %7) local_unnamed_addr #0 !dbg !4 {\n %9 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %10 = shl i32 %9, 3, !dbg !13\n %11 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !14\n %12 = and i32 %11, 7, !dbg !14\n %13 = or disjoint i32 %10, %12, !dbg !15\n %14 = icmp slt i32 %13, %4, !dbg !16\n %15 = sext i32 %13 to i64, !dbg !17\n %16 = getelementptr float, ptr addrspace(1) %0, i64 %15, !dbg !17\n %17 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %16, i1 %14) #2, !dbg !18\n %18 = bitcast i32 %17 to float, !dbg !18\n %19 = getelementptr float, ptr addrspace(1) %1, i64 %15, !dbg !19\n %20 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %19, i1 %14) #2, !dbg !20\n %21 = bitcast i32 %20 to float, !dbg !20\n %22 = getelementptr float, ptr addrspace(1) %2, i64 %15, !dbg !21\n %23 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %22, i1 %14) #2, !dbg !22\n %24 = bitcast i32 %23 to float, !dbg !22\n %25 = fmul float %18, %21, !dbg !23\n %26 = fmul float %5, %25, !dbg !24\n %27 = fadd float %26, %24, !dbg !25\n %28 = getelementptr float, ptr addrspace(1) %3, i64 %15, !dbg !26\n %29 = and i32 %11, 120, !dbg !27\n %30 = icmp eq i32 %29, 0, !dbg !27\n %31 = bitcast float %27 to i32, !dbg !27\n %32 = and i1 %30, %14, !dbg !27\n tail call void asm sideeffect \"@$2 st.global.b32 [ $1 + 0 ], { $0 };\", \"r,l,b\"(i32 %31, ptr addrspace(1) %28, i1 %32) #2, !dbg !27\n ret void, !dbg !28\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"128\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"fused_op_kernel\", linkageName: \"fused_op_kernel\", scope: !1, file: !1, line: 347, type: !5, scopeLine: 347, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !7, !9, !8, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"float\", size: 32, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 357, column: 24, scope: !4)\n!13 = !DILocation(line: 358, column: 20, scope: !4)\n!14 = !DILocation(line: 358, column: 46, scope: !4)\n!15 = !DILocation(line: 358, column: 33, scope: !4)\n!16 = !DILocation(line: 359, column: 21, scope: !4)\n!17 = !DILocation(line: 361, column: 24, scope: !4)\n!18 = !DILocation(line: 361, column: 16, scope: !4)\n!19 = !DILocation(line: 362, column: 24, scope: !4)\n!20 = !DILocation(line: 362, column: 16, scope: !4)\n!21 = !DILocation(line: 363, column: 24, scope: !4)\n!22 = !DILocation(line: 363, column: 16, scope: !4)\n!23 = !DILocation(line: 365, column: 17, scope: !4)\n!24 = !DILocation(line: 365, column: 21, scope: !4)\n!25 = !DILocation(line: 365, column: 36, scope: !4)\n!26 = !DILocation(line: 369, column: 26, scope: !4)\n!27 = !DILocation(line: 369, column: 35, scope: !4)\n!28 = !DILocation(line: 369, column: 4, scope: !4)\n","fused_op_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.7\n.target sm_90a\n.address_size 64\n\n\t// .globl\tfused_op_kernel // -- Begin function fused_op_kernel\n // @fused_op_kernel\n.visible .entry fused_op_kernel(\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_2,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_3,\n\t.param .u32 fused_op_kernel_param_4,\n\t.param .f32 fused_op_kernel_param_5,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_6,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_7\n)\n.reqntid 128\n{\n\t.reg .pred \t%p<4>;\n\t.reg .b32 \t%r<14>;\n\t.reg .b64 \t%rd<10>;\n\t.loc\t1 347 0 // generate_examples.py:347:0\n$L__func_begin0:\n\t.loc\t1 347 0 // generate_examples.py:347:0\n\n// %bb.0:\n\tld.param.b64 \t%rd5, [fused_op_kernel_param_0];\n\tld.param.b64 \t%rd6, [fused_op_kernel_param_1];\n$L__tmp0:\n\t.loc\t1 357 24 // generate_examples.py:357:24\n\tmov.u32 \t%r5, %ctaid.x;\n\t.loc\t1 358 20 // generate_examples.py:358:20\n\tshl.b32 \t%r6, %r5, 3;\n\tld.param.b64 \t%rd7, [fused_op_kernel_param_2];\n\tld.param.b64 \t%rd8, [fused_op_kernel_param_3];\n\t.loc\t1 358 46 // generate_examples.py:358:46\n\tmov.u32 \t%r7, %tid.x;\n\tand.b32 \t%r8, %r7, 7;\n\tld.param.b32 \t%r9, [fused_op_kernel_param_4];\n\t.loc\t1 358 33 // generate_examples.py:358:33\n\tor.b32 \t%r10, %r6, %r8;\n\tld.param.b32 \t%r11, [fused_op_kernel_param_5];\n\t.loc\t1 359 21 // generate_examples.py:359:21\n\tsetp.lt.s32 \t%p1, %r10, %r9;\n\t.loc\t1 361 24 // generate_examples.py:361:24\n\tmul.wide.s32 \t%rd9, %r10, 4;\n\tadd.s64 \t%rd1, %rd5, %rd9;\n\t.loc\t1 361 16 // generate_examples.py:361:16\n\t// begin inline asm\n\tmov.u32 %r1, 0x0;\n\t@%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];\n\t// end inline asm\n\t.loc\t1 362 24 // generate_examples.py:362:24\n\tadd.s64 \t%rd2, %rd6, %rd9;\n\t.loc\t1 362 16 // generate_examples.py:362:16\n\t// begin inline asm\n\tmov.u32 %r2, 0x0;\n\t@%p1 ld.global.b32 { %r2 }, [ %rd2 + 0 ];\n\t// end inline asm\n\t.loc\t1 363 24 // generate_examples.py:363:24\n\tadd.s64 \t%rd3, %rd7, %rd9;\n\t.loc\t1 363 16 // generate_examples.py:363:16\n\t// begin inline asm\n\tmov.u32 %r3, 0x0;\n\t@%p1 ld.global.b32 { %r3 }, [ %rd3 + 0 ];\n\t// end inline asm\n\t.loc\t1 365 17 // generate_examples.py:365:17\n\tmul.f32 \t%r12, %r1, %r2;\n\t.loc\t1 365 36 // generate_examples.py:365:36\n\tfma.rn.f32 \t%r4, %r11, %r12, %r3;\n\t.loc\t1 369 26 // generate_examples.py:369:26\n\tadd.s64 \t%rd4, %rd8, %rd9;\n\t.loc\t1 369 35 // generate_examples.py:369:35\n\tand.b32 \t%r13, %r7, 120;\n\tsetp.eq.b32 \t%p3, %r13, 0;\n\tand.pred \t%p2, %p3, %p1;\n\t// begin inline asm\n\t@%p2 st.global.b32 [ %rd4 + 0 ], { %r4 };\n\t// end inline asm\n\t.loc\t1 369 4 // generate_examples.py:369:4\n\tret;\n$L__tmp1:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 0 // DW_CHILDREN_no\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 131 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0x7c DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 45\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 45\n.b8 114\n.b8 101\n.b8 103\n.b8 101\n.b8 110\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","fused_op_kernel.json":"{\"hash\": \"1aa7bd5d4058894489843d4e8c2bfc6c480457536a3cc8e94bc4553e36fb837e\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 4, \"num_ctas\": 1, \"num_stages\": 3, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.7.0\", \"tensordesc_meta\": [], \"shared\": 0, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"fused_op_kernel\"}","fused_op_kernel.sass":"Function:fused_op_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 347\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 41\n /*0010*/ S2R R12, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*0020*/ LDC.64 R2, c[0x0][0x210] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*0030*/ ULDC UR6, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 62\n /*0040*/ MOV R13, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 357\n\t//## File \".nv_debug_ptx_txt\", line 35\n /*0050*/ S2R R11, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 58\n /*0060*/ LDC.64 R4, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0070*/ LDC.64 R6, c[0x0][0x220] ;\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0080*/ HFMA2.MMA R10, -RZ, RZ, 0, 0 ;\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0090*/ ULDC.64 UR4, c[0x0][0x208] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*00a0*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*00b0*/ LOP3.LUT R0, R12, 0x7, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*00c0*/ LEA R11, R11, R0, 0x3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*00d0*/ HFMA2.MMA R0, -RZ, RZ, 0, 0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*00e0*/ ISETP.GE.AND P0, PT, R11.reuse, UR6, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*00f0*/ IMAD.WIDE R2, R11, 0x4, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 58\n /*0100*/ IMAD.WIDE R4, R11, 0x4, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0110*/ IMAD.WIDE R6, R11, 0x4, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*0120*/ @!P0 LDG.E R0, desc[UR4][R2.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 62\n /*0130*/ @!P0 LDG.E R13, desc[UR4][R4.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0140*/ @!P0 LDG.E R10, desc[UR4][R6.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0150*/ LOP3.LUT P0, RZ, R12, 0x78, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0160*/ IMAD.WIDE R8, R11, 0x4, R8 ;\n /*0170*/ ISETP.LT.AND P0, PT, R11, UR6, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 365\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0180*/ ULDC UR6, c[0x0][0x234] ;\n\t//## File \".nv_debug_ptx_txt\", line 72\n /*0190*/ FMUL R13, R13, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*01a0*/ FFMA R13, R13, UR6, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 82\n /*01b0*/ @!P0 EXIT ;\n /*01c0*/ STG.E desc[UR4][R8.64], R13 ;\n\t//## File \".nv_debug_ptx_txt\", line 85\n /*01d0*/ EXIT ;\n.L_x_0:\n /*01e0*/ BRA `(.L_x_0);\n /*01f0*/ NOP;\n /*0200*/ NOP;\n /*0210*/ NOP;\n /*0220*/ NOP;\n /*0230*/ NOP;\n /*0240*/ NOP;\n /*0250*/ NOP;\n /*0260*/ NOP;\n /*0270*/ NOP;\n.L_x_1:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":346,"end_line":369,"code":" @triton.jit\n def fused_op_kernel(\n a_ptr,\n b_ptr,\n c_ptr,\n output_ptr,\n n_elements,\n scale_factor: float,\n ACTIVATION: tl.constexpr,\n BLOCK_SIZE: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n mask = offsets < n_elements\n\n a = tl.load(a_ptr + offsets, mask=mask)\n b = tl.load(b_ptr + offsets, mask=mask)\n c = tl.load(c_ptr + offsets, mask=mask)\n\n result = a * b * scale_factor + c\n if ACTIVATION == \"relu\":\n result = tl.where(result > 0, result, 0.0)\n\n tl.store(output_ptr + offsets, result, mask=mask)\n"}}} -{"event_type":"compilation","pid":1724463,"stack":[{"line":198,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":431,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op(x, y, z, scale_factor=1.0, activation=\"relu\").sum()"},{"line":404,"name":"fused_op","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op_kernel[grid]("},{"line":370,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":739,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":884,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-21T21:13:23.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"64906c37b60403702b7090fe8d7fb463d9694159865d6fa84065c0c42669be02","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"fused_op_kernel","num_ctas":1,"num_stages":3,"num_warps":4,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":0,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.7.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[],"(6,)":[],"(7,)":[]},"src_cache_key":"f2e5ae6c41d3bfc26af60eada8fff38a0b318cffe77df20b32d117346110ff07","src_constants":{"(6,)":"relu","(7,)":8},"times":{"ir_initialization":1263,"lowering_stages":[],"store_results":0}},"file_path":{"fused_op_kernel.source":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.source","fused_op_kernel.ttir":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.ttir","fused_op_kernel.ttgir":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.ttgir","fused_op_kernel.llir":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.llir","fused_op_kernel.ptx":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.ptx","fused_op_kernel.cubin":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.cubin","fused_op_kernel.json":"/home/user/.triton/cache/MSIGYN5WAQBXAK3QSD7I275UMPMWSQKZQZOW7KCAMXAMIJTJXYBA/fused_op_kernel.json"},"file_content":{"fused_op_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:0)\n#loc21 = loc(\"a_ptr\"(#loc))\n#loc22 = loc(\"b_ptr\"(#loc))\n#loc23 = loc(\"c_ptr\"(#loc))\n#loc24 = loc(\"output_ptr\"(#loc))\n#loc25 = loc(\"n_elements\"(#loc))\n#loc26 = loc(\"scale_factor\"(#loc))\nmodule {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %cst = arith.constant dense<0.000000e+00> : tensor<8xf32> loc(#loc1)\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc27)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc28)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc29)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32> loc(#loc30)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32> loc(#loc30)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32> loc(#loc31)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32> loc(#loc31)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc32)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc32)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr> loc(#loc33)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc34)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc34)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr> loc(#loc35)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc36)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc36)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr> loc(#loc37)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32> loc(#loc38)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32> loc(#loc39)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32> loc(#loc39)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32> loc(#loc40)\n %result_13 = arith.cmpf ogt, %result_12, %cst : tensor<8xf32> loc(#loc41)\n %result_14 = arith.select %result_13, %result_12, %cst : tensor<8xi1>, tensor<8xf32> loc(#loc42)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc18)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc18)\n tt.store %1, %result_14, %mask_3 : tensor<8x!tt.ptr> loc(#loc19)\n tt.return loc(#loc20)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:24)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:20)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:46)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:33)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:24)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:16)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:24)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:16)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:24)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:16)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:17)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:21)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:36)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:35)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:46)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:26)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:35)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:4)\n#loc27 = loc(\"pid\"(#loc2))\n#loc28 = loc(\"offsets\"(#loc3))\n#loc29 = loc(\"offsets\"(#loc4))\n#loc30 = loc(\"offsets\"(#loc5))\n#loc31 = loc(\"mask\"(#loc6))\n#loc32 = loc(\"a\"(#loc7))\n#loc33 = loc(\"a\"(#loc8))\n#loc34 = loc(\"b\"(#loc9))\n#loc35 = loc(\"b\"(#loc10))\n#loc36 = loc(\"c\"(#loc11))\n#loc37 = loc(\"c\"(#loc12))\n#loc38 = loc(\"result\"(#loc13))\n#loc39 = loc(\"result\"(#loc14))\n#loc40 = loc(\"result\"(#loc15))\n#loc41 = loc(\"result\"(#loc16))\n#loc42 = loc(\"result\"(#loc17))\n","fused_op_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:0)\n#loc21 = loc(\"a_ptr\"(#loc))\n#loc22 = loc(\"b_ptr\"(#loc))\n#loc23 = loc(\"c_ptr\"(#loc))\n#loc24 = loc(\"output_ptr\"(#loc))\n#loc25 = loc(\"n_elements\"(#loc))\n#loc26 = loc(\"scale_factor\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 4 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %cst = arith.constant dense<0.000000e+00> : tensor<8xf32, #blocked> loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc27)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc28)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #blocked> loc(#loc29)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32, #blocked> loc(#loc30)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32, #blocked> loc(#loc30)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32, #blocked> loc(#loc31)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32, #blocked> loc(#loc31)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc32)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc32)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc33)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc34)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc34)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc35)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc36)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc36)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc37)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32, #blocked> loc(#loc38)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32, #blocked> loc(#loc39)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32, #blocked> loc(#loc39)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32, #blocked> loc(#loc40)\n %result_13 = arith.cmpf ogt, %result_12, %cst : tensor<8xf32, #blocked> loc(#loc41)\n %result_14 = arith.select %result_13, %result_12, %cst : tensor<8xi1, #blocked>, tensor<8xf32, #blocked> loc(#loc42)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc18)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc18)\n tt.store %1, %result_14, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc19)\n tt.return loc(#loc20)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:24)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:20)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:46)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:33)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:24)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:16)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:24)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:16)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:24)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:16)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:17)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:21)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:36)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:35)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:46)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:26)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:35)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:4)\n#loc27 = loc(\"pid\"(#loc2))\n#loc28 = loc(\"offsets\"(#loc3))\n#loc29 = loc(\"offsets\"(#loc4))\n#loc30 = loc(\"offsets\"(#loc5))\n#loc31 = loc(\"mask\"(#loc6))\n#loc32 = loc(\"a\"(#loc7))\n#loc33 = loc(\"a\"(#loc8))\n#loc34 = loc(\"b\"(#loc9))\n#loc35 = loc(\"b\"(#loc10))\n#loc36 = loc(\"c\"(#loc11))\n#loc37 = loc(\"c\"(#loc12))\n#loc38 = loc(\"result\"(#loc13))\n#loc39 = loc(\"result\"(#loc14))\n#loc40 = loc(\"result\"(#loc15))\n#loc41 = loc(\"result\"(#loc16))\n#loc42 = loc(\"result\"(#loc17))\n","fused_op_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @fused_op_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, i32 %4, float %5, ptr addrspace(1) readnone captures(none) %6, ptr addrspace(1) readnone captures(none) %7) local_unnamed_addr #0 !dbg !4 {\n %9 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %10 = shl i32 %9, 3, !dbg !13\n %11 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !14\n %12 = and i32 %11, 7, !dbg !14\n %13 = or disjoint i32 %10, %12, !dbg !15\n %14 = icmp slt i32 %13, %4, !dbg !16\n %15 = sext i32 %13 to i64, !dbg !17\n %16 = getelementptr float, ptr addrspace(1) %0, i64 %15, !dbg !17\n %17 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %16, i1 %14) #2, !dbg !18\n %18 = bitcast i32 %17 to float, !dbg !18\n %19 = getelementptr float, ptr addrspace(1) %1, i64 %15, !dbg !19\n %20 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %19, i1 %14) #2, !dbg !20\n %21 = bitcast i32 %20 to float, !dbg !20\n %22 = getelementptr float, ptr addrspace(1) %2, i64 %15, !dbg !21\n %23 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %22, i1 %14) #2, !dbg !22\n %24 = bitcast i32 %23 to float, !dbg !22\n %25 = fmul float %18, %21, !dbg !23\n %26 = fmul float %5, %25, !dbg !24\n %27 = fadd float %26, %24, !dbg !25\n %28 = fcmp ogt float %27, 0.000000e+00, !dbg !26\n %29 = select i1 %28, float %27, float 0.000000e+00, !dbg !27\n %30 = getelementptr float, ptr addrspace(1) %3, i64 %15, !dbg !28\n %31 = and i32 %11, 120, !dbg !29\n %32 = icmp eq i32 %31, 0, !dbg !29\n %33 = bitcast float %29 to i32, !dbg !29\n %34 = and i1 %32, %14, !dbg !29\n tail call void asm sideeffect \"@$2 st.global.b32 [ $1 + 0 ], { $0 };\", \"r,l,b\"(i32 %33, ptr addrspace(1) %30, i1 %34) #2, !dbg !29\n ret void, !dbg !30\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"128\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"fused_op_kernel\", linkageName: \"fused_op_kernel\", scope: !1, file: !1, line: 347, type: !5, scopeLine: 347, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !7, !9, !8, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"float\", size: 32, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 357, column: 24, scope: !4)\n!13 = !DILocation(line: 358, column: 20, scope: !4)\n!14 = !DILocation(line: 358, column: 46, scope: !4)\n!15 = !DILocation(line: 358, column: 33, scope: !4)\n!16 = !DILocation(line: 359, column: 21, scope: !4)\n!17 = !DILocation(line: 361, column: 24, scope: !4)\n!18 = !DILocation(line: 361, column: 16, scope: !4)\n!19 = !DILocation(line: 362, column: 24, scope: !4)\n!20 = !DILocation(line: 362, column: 16, scope: !4)\n!21 = !DILocation(line: 363, column: 24, scope: !4)\n!22 = !DILocation(line: 363, column: 16, scope: !4)\n!23 = !DILocation(line: 365, column: 17, scope: !4)\n!24 = !DILocation(line: 365, column: 21, scope: !4)\n!25 = !DILocation(line: 365, column: 36, scope: !4)\n!26 = !DILocation(line: 367, column: 35, scope: !4)\n!27 = !DILocation(line: 367, column: 46, scope: !4)\n!28 = !DILocation(line: 369, column: 26, scope: !4)\n!29 = !DILocation(line: 369, column: 35, scope: !4)\n!30 = !DILocation(line: 369, column: 4, scope: !4)\n","fused_op_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.7\n.target sm_90a\n.address_size 64\n\n\t// .globl\tfused_op_kernel // -- Begin function fused_op_kernel\n // @fused_op_kernel\n.visible .entry fused_op_kernel(\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_2,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_3,\n\t.param .u32 fused_op_kernel_param_4,\n\t.param .f32 fused_op_kernel_param_5,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_6,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_7\n)\n.reqntid 128\n{\n\t.reg .pred \t%p<4>;\n\t.reg .b32 \t%r<15>;\n\t.reg .b64 \t%rd<10>;\n\t.loc\t1 347 0 // generate_examples.py:347:0\n$L__func_begin0:\n\t.loc\t1 347 0 // generate_examples.py:347:0\n\n// %bb.0:\n\tld.param.b64 \t%rd5, [fused_op_kernel_param_0];\n\tld.param.b64 \t%rd6, [fused_op_kernel_param_1];\n$L__tmp0:\n\t.loc\t1 357 24 // generate_examples.py:357:24\n\tmov.u32 \t%r5, %ctaid.x;\n\t.loc\t1 358 20 // generate_examples.py:358:20\n\tshl.b32 \t%r6, %r5, 3;\n\tld.param.b64 \t%rd7, [fused_op_kernel_param_2];\n\tld.param.b64 \t%rd8, [fused_op_kernel_param_3];\n\t.loc\t1 358 46 // generate_examples.py:358:46\n\tmov.u32 \t%r7, %tid.x;\n\tand.b32 \t%r8, %r7, 7;\n\tld.param.b32 \t%r9, [fused_op_kernel_param_4];\n\t.loc\t1 358 33 // generate_examples.py:358:33\n\tor.b32 \t%r10, %r6, %r8;\n\tld.param.b32 \t%r11, [fused_op_kernel_param_5];\n\t.loc\t1 359 21 // generate_examples.py:359:21\n\tsetp.lt.s32 \t%p1, %r10, %r9;\n\t.loc\t1 361 24 // generate_examples.py:361:24\n\tmul.wide.s32 \t%rd9, %r10, 4;\n\tadd.s64 \t%rd1, %rd5, %rd9;\n\t.loc\t1 361 16 // generate_examples.py:361:16\n\t// begin inline asm\n\tmov.u32 %r1, 0x0;\n\t@%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];\n\t// end inline asm\n\t.loc\t1 362 24 // generate_examples.py:362:24\n\tadd.s64 \t%rd2, %rd6, %rd9;\n\t.loc\t1 362 16 // generate_examples.py:362:16\n\t// begin inline asm\n\tmov.u32 %r2, 0x0;\n\t@%p1 ld.global.b32 { %r2 }, [ %rd2 + 0 ];\n\t// end inline asm\n\t.loc\t1 363 24 // generate_examples.py:363:24\n\tadd.s64 \t%rd3, %rd7, %rd9;\n\t.loc\t1 363 16 // generate_examples.py:363:16\n\t// begin inline asm\n\tmov.u32 %r3, 0x0;\n\t@%p1 ld.global.b32 { %r3 }, [ %rd3 + 0 ];\n\t// end inline asm\n\t.loc\t1 365 17 // generate_examples.py:365:17\n\tmul.f32 \t%r12, %r1, %r2;\n\t.loc\t1 365 36 // generate_examples.py:365:36\n\tfma.rn.f32 \t%r13, %r11, %r12, %r3;\n\t.loc\t1 367 46 // generate_examples.py:367:46\n\tmax.f32 \t%r4, %r13, 0f00000000;\n\t.loc\t1 369 26 // generate_examples.py:369:26\n\tadd.s64 \t%rd4, %rd8, %rd9;\n\t.loc\t1 369 35 // generate_examples.py:369:35\n\tand.b32 \t%r14, %r7, 120;\n\tsetp.eq.b32 \t%p3, %r14, 0;\n\tand.pred \t%p2, %p3, %p1;\n\t// begin inline asm\n\t@%p2 st.global.b32 [ %rd4 + 0 ], { %r4 };\n\t// end inline asm\n\t.loc\t1 369 4 // generate_examples.py:369:4\n\tret;\n$L__tmp1:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 0 // DW_CHILDREN_no\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 131 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0x7c DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 45\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 45\n.b8 114\n.b8 101\n.b8 103\n.b8 101\n.b8 110\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","fused_op_kernel.json":"{\"hash\": \"64906c37b60403702b7090fe8d7fb463d9694159865d6fa84065c0c42669be02\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 4, \"num_ctas\": 1, \"num_stages\": 3, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.7.0\", \"tensordesc_meta\": [], \"shared\": 0, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"fused_op_kernel\"}","fused_op_kernel.sass":"Function:fused_op_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 347\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 41\n /*0010*/ S2R R12, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*0020*/ LDC.64 R2, c[0x0][0x210] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*0030*/ ULDC UR6, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 62\n /*0040*/ MOV R13, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 357\n\t//## File \".nv_debug_ptx_txt\", line 35\n /*0050*/ S2R R11, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 58\n /*0060*/ LDC.64 R4, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0070*/ LDC.64 R6, c[0x0][0x220] ;\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0080*/ HFMA2.MMA R10, -RZ, RZ, 0, 0 ;\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0090*/ ULDC.64 UR4, c[0x0][0x208] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*00a0*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*00b0*/ LOP3.LUT R0, R12, 0x7, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*00c0*/ LEA R11, R11, R0, 0x3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*00d0*/ HFMA2.MMA R0, -RZ, RZ, 0, 0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*00e0*/ ISETP.GE.AND P0, PT, R11.reuse, UR6, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*00f0*/ IMAD.WIDE R2, R11, 0x4, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 58\n /*0100*/ IMAD.WIDE R4, R11, 0x4, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0110*/ IMAD.WIDE R6, R11, 0x4, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*0120*/ @!P0 LDG.E R0, desc[UR4][R2.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 62\n /*0130*/ @!P0 LDG.E R13, desc[UR4][R4.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0140*/ @!P0 LDG.E R10, desc[UR4][R6.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*0150*/ LOP3.LUT P0, RZ, R12, 0x78, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*0160*/ IMAD.WIDE R2, R11, 0x4, R8 ;\n /*0170*/ ISETP.LT.AND P0, PT, R11, UR6, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 365\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0180*/ ULDC UR6, c[0x0][0x234] ;\n\t//## File \".nv_debug_ptx_txt\", line 72\n /*0190*/ FMUL R13, R13, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*01a0*/ FFMA R10, R13, UR6, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 367\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*01b0*/ FMNMX R9, RZ, R10, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 84\n /*01c0*/ @!P0 EXIT ;\n /*01d0*/ STG.E desc[UR4][R2.64], R9 ;\n\t//## File \".nv_debug_ptx_txt\", line 87\n /*01e0*/ EXIT ;\n.L_x_0:\n /*01f0*/ BRA `(.L_x_0);\n /*0200*/ NOP;\n /*0210*/ NOP;\n /*0220*/ NOP;\n /*0230*/ NOP;\n /*0240*/ NOP;\n /*0250*/ NOP;\n /*0260*/ NOP;\n /*0270*/ NOP;\n.L_x_1:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":346,"end_line":369,"code":" @triton.jit\n def fused_op_kernel(\n a_ptr,\n b_ptr,\n c_ptr,\n output_ptr,\n n_elements,\n scale_factor: float,\n ACTIVATION: tl.constexpr,\n BLOCK_SIZE: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n mask = offsets < n_elements\n\n a = tl.load(a_ptr + offsets, mask=mask)\n b = tl.load(b_ptr + offsets, mask=mask)\n c = tl.load(c_ptr + offsets, mask=mask)\n\n result = a * b * scale_factor + c\n if ACTIVATION == \"relu\":\n result = tl.where(result > 0, result, 0.0)\n\n tl.store(output_ptr + offsets, result, mask=mask)\n"}}} +{"event_type":"compilation","pid":4104992,"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":272,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":757,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":902,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[["tt.divisibility",16]],"(5,)":[["tt.divisibility",16]],"(6,)":[["tt.divisibility",16]],"(7,)":[],"(8,)":[["tt.divisibility",16]],"(9,)":[],"(10,)":[["tt.divisibility",16]],"(11,)":[],"(12,)":[],"(13,)":[],"(14,)":[],"(15,)":[],"(16,)":[]},"src_cache_key":"528f8911847b1f39f462157322ec9a674a5565b9b74d621d5ff09625eb710f38","src_constants":{"(7,)":1,"(9,)":1,"(11,)":1,"(12,)":16,"(13,)":16,"(14,)":16,"(15,)":1,"(16,)":null},"times":{"ir_initialization":578768,"lowering_stages":[],"store_results":0}},"file_path":{"matmul_kernel.source":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.source","matmul_kernel.ttir":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.ttir","matmul_kernel.ttgir":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.ttgir","matmul_kernel.llir":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.llir","matmul_kernel.ptx":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.ptx","matmul_kernel.cubin":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.cubin","matmul_kernel.json":"/home/user/.triton/cache/FNNRJHU5LLSRMH4YTBEIFLIRCDN72CXDAANQ6EB3VLT6YZBHIZGA/matmul_kernel.json"},"file_content":{"matmul_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:5)\n#loc46 = loc(\"a\"(#loc))\n#loc47 = loc(\"b\"(#loc))\n#loc48 = loc(\"c\"(#loc))\n#loc49 = loc(\"M\"(#loc))\n#loc50 = loc(\"N\"(#loc))\n#loc51 = loc(\"K\"(#loc))\n#loc52 = loc(\"stride_am\"(#loc))\n#loc53 = loc(\"stride_bk\"(#loc))\n#loc54 = loc(\"stride_cm\"(#loc))\nmodule {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %cst = arith.constant dense<0.000000e+00> : tensor<16x16xf16> loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc2)\n %cst_0 = arith.constant dense<16> : tensor<16x16xi32> loc(#loc1)\n %cst_1 = arith.constant dense<0.000000e+00> : tensor<16x16xf32> loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc55)\n %num_pid_m = arith.addi %M, %c15_i32 : i32 loc(#loc97)\n %num_pid_m_2 = arith.divsi %num_pid_m, %c16_i32 : i32 loc(#loc98)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc99)\n %num_pid_n_3 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc100)\n %group_id = arith.divsi %pid, %num_pid_n_3 : i32 loc(#loc58)\n %group_size = arith.subi %num_pid_m_2, %group_id : i32 loc(#loc59)\n %group_size_4 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc60)\n %pid_m = arith.remsi %pid, %group_size_4 : i32 loc(#loc61)\n %pid_m_5 = arith.addi %group_id, %pid_m : i32 loc(#loc62)\n %pid_n = arith.remsi %pid, %num_pid_n_3 : i32 loc(#loc63)\n %pid_n_6 = arith.divsi %pid_n, %group_size_4 : i32 loc(#loc64)\n %offs_am = arith.muli %pid_m_5, %c16_i32 : i32 loc(#loc65)\n %offs_am_7 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc66)\n %offs_am_8 = tt.splat %offs_am : i32 -> tensor<16xi32> loc(#loc65)\n %offs_am_9 = arith.addi %offs_am_8, %offs_am_7 : tensor<16xi32> loc(#loc65)\n %offs_am_10 = tt.splat %M : i32 -> tensor<16xi32> loc(#loc67)\n %offs_am_11 = arith.remsi %offs_am_9, %offs_am_10 : tensor<16xi32> loc(#loc67)\n %offs_bn = arith.muli %pid_n_6, %c16_i32 : i32 loc(#loc68)\n %offs_bn_12 = tt.splat %offs_bn : i32 -> tensor<16xi32> loc(#loc68)\n %offs_bn_13 = arith.addi %offs_bn_12, %offs_am_7 : tensor<16xi32> loc(#loc68)\n %offs_bn_14 = tt.splat %N : i32 -> tensor<16xi32> loc(#loc69)\n %offs_bn_15 = arith.remsi %offs_bn_13, %offs_bn_14 : tensor<16xi32> loc(#loc69)\n %a_ptrs = tt.expand_dims %offs_am_11 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc70)\n %a_ptrs_16 = tt.splat %stride_am : i32 -> tensor<16x1xi32> loc(#loc70)\n %a_ptrs_17 = arith.muli %a_ptrs, %a_ptrs_16 : tensor<16x1xi32> loc(#loc70)\n %a_ptrs_18 = tt.expand_dims %offs_am_7 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc71)\n %a_ptrs_19 = tt.broadcast %a_ptrs_17 : tensor<16x1xi32> -> tensor<16x16xi32> loc(#loc70)\n %a_ptrs_20 = tt.broadcast %a_ptrs_18 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc70)\n %a_ptrs_21 = arith.addi %a_ptrs_19, %a_ptrs_20 : tensor<16x16xi32> loc(#loc70)\n %a_ptrs_22 = tt.splat %a : !tt.ptr -> tensor<16x16x!tt.ptr> loc(#loc72)\n %a_ptrs_23 = tt.addptr %a_ptrs_22, %a_ptrs_21 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc72)\n %b_ptrs = tt.expand_dims %offs_am_7 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc73)\n %b_ptrs_24 = tt.splat %stride_bk : i32 -> tensor<16x1xi32> loc(#loc73)\n %b_ptrs_25 = arith.muli %b_ptrs, %b_ptrs_24 : tensor<16x1xi32> loc(#loc73)\n %b_ptrs_26 = tt.expand_dims %offs_bn_15 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc74)\n %b_ptrs_27 = tt.broadcast %b_ptrs_25 : tensor<16x1xi32> -> tensor<16x16xi32> loc(#loc73)\n %b_ptrs_28 = tt.broadcast %b_ptrs_26 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc73)\n %b_ptrs_29 = arith.addi %b_ptrs_27, %b_ptrs_28 : tensor<16x16xi32> loc(#loc73)\n %b_ptrs_30 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr> loc(#loc75)\n %b_ptrs_31 = tt.addptr %b_ptrs_30, %b_ptrs_29 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc75)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc76)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc77)\n %accumulator:3 = scf.for %k = %c0_i32 to %1 step %c1_i32 iter_args(%a_ptrs_46 = %a_ptrs_23, %b_ptrs_47 = %b_ptrs_31, %accumulator_48 = %cst_1) -> (tensor<16x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<16x16xf32>) : i32 {\n %a_block = arith.muli %k, %c16_i32 : i32 loc(#loc79)\n %a_block_49 = arith.subi %K, %a_block : i32 loc(#loc80)\n %a_block_50 = tt.splat %a_block_49 : i32 -> tensor<1x16xi32> loc(#loc81)\n %a_block_51 = arith.cmpi slt, %a_ptrs_18, %a_block_50 : tensor<1x16xi32> loc(#loc81)\n %a_block_52 = tt.broadcast %a_block_51 : tensor<1x16xi1> -> tensor<16x16xi1> loc(#loc82)\n %a_block_53 = tt.load %a_ptrs_46, %a_block_52, %cst : tensor<16x16x!tt.ptr> loc(#loc82)\n %b_block = tt.splat %a_block_49 : i32 -> tensor<16x1xi32> loc(#loc83)\n %b_block_54 = arith.cmpi slt, %b_ptrs, %b_block : tensor<16x1xi32> loc(#loc83)\n %b_block_55 = tt.broadcast %b_block_54 : tensor<16x1xi1> -> tensor<16x16xi1> loc(#loc84)\n %b_block_56 = tt.load %b_ptrs_47, %b_block_55, %cst : tensor<16x16x!tt.ptr> loc(#loc84)\n %accumulator_57 = tt.dot %a_block_53, %b_block_56, %accumulator_48, inputPrecision = tf32 : tensor<16x16xf16> * tensor<16x16xf16> -> tensor<16x16xf32> loc(#loc85)\n %a_ptrs_58 = tt.addptr %a_ptrs_46, %cst_0 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc86)\n %b_ptrs_59 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc87)\n %b_ptrs_60 = tt.splat %b_ptrs_59 : i32 -> tensor<16x16xi32> loc(#loc88)\n %b_ptrs_61 = tt.addptr %b_ptrs_47, %b_ptrs_60 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc88)\n scf.yield %a_ptrs_58, %b_ptrs_61, %accumulator_57 : tensor<16x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<16x16xf32> loc(#loc2)\n } loc(#loc102)\n %c_block = arith.truncf %accumulator#2 : tensor<16x16xf32> to tensor<16x16xf16> loc(#loc89)\n %c_ptrs = tt.expand_dims %offs_am_9 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc90)\n %c_ptrs_32 = tt.splat %stride_cm : i32 -> tensor<16x1xi32> loc(#loc91)\n %c_ptrs_33 = arith.muli %c_ptrs_32, %c_ptrs : tensor<16x1xi32> loc(#loc91)\n %c_ptrs_34 = tt.splat %c : !tt.ptr -> tensor<16x1x!tt.ptr> loc(#loc92)\n %c_ptrs_35 = tt.addptr %c_ptrs_34, %c_ptrs_33 : tensor<16x1x!tt.ptr>, tensor<16x1xi32> loc(#loc92)\n %c_ptrs_36 = tt.expand_dims %offs_bn_13 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc93)\n %c_ptrs_37 = tt.broadcast %c_ptrs_35 : tensor<16x1x!tt.ptr> -> tensor<16x16x!tt.ptr> loc(#loc92)\n %c_ptrs_38 = tt.broadcast %c_ptrs_36 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc92)\n %c_ptrs_39 = tt.addptr %c_ptrs_37, %c_ptrs_38 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc92)\n %c_mask = tt.splat %M : i32 -> tensor<16x1xi32> loc(#loc94)\n %c_mask_40 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<16x1xi32> loc(#loc94)\n %c_mask_41 = tt.splat %N : i32 -> tensor<1x16xi32> loc(#loc95)\n %c_mask_42 = arith.cmpi slt, %c_ptrs_36, %c_mask_41 : tensor<1x16xi32> loc(#loc95)\n %c_mask_43 = tt.broadcast %c_mask_40 : tensor<16x1xi1> -> tensor<16x16xi1> loc(#loc96)\n %c_mask_44 = tt.broadcast %c_mask_42 : tensor<1x16xi1> -> tensor<16x16xi1> loc(#loc96)\n %c_mask_45 = arith.andi %c_mask_43, %c_mask_44 : tensor<16x16xi1> loc(#loc96)\n tt.store %c_ptrs_39, %c_block, %c_mask_45 : tensor<16x16x!tt.ptr> loc(#loc45)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:9)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:15)\n#loc4 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:13)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:21)\n#loc6 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:12)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:21)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:20)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:26)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:22)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:32)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:17)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:18)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:17)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:20)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:43)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:19)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:20)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:19)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:23)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:54)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:23)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:53)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:27)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:52)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:48)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:30)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":329:23)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:30)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":332:23)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:28)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:13)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:23)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:13)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:19)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:34)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:22)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:18)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:65)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:19)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:44)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:18)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:9)\n#loc55 = loc(\"pid\"(#loc3))\n#loc56 = loc(\"num_pid_m\"(#loc5))\n#loc57 = loc(\"num_pid_n\"(#loc7))\n#loc58 = loc(\"group_id\"(#loc8))\n#loc59 = loc(\"group_size\"(#loc9))\n#loc60 = loc(\"group_size\"(#loc10))\n#loc61 = loc(\"pid_m\"(#loc11))\n#loc62 = loc(\"pid_m\"(#loc12))\n#loc63 = loc(\"pid_n\"(#loc13))\n#loc64 = loc(\"pid_n\"(#loc14))\n#loc65 = loc(\"offs_am\"(#loc15))\n#loc66 = loc(\"offs_am\"(#loc16))\n#loc67 = loc(\"offs_am\"(#loc17))\n#loc68 = loc(\"offs_bn\"(#loc18))\n#loc69 = loc(\"offs_bn\"(#loc19))\n#loc70 = loc(\"a_ptrs\"(#loc20))\n#loc71 = loc(\"a_ptrs\"(#loc21))\n#loc72 = loc(\"a_ptrs\"(#loc22))\n#loc73 = loc(\"b_ptrs\"(#loc23))\n#loc74 = loc(\"b_ptrs\"(#loc24))\n#loc75 = loc(\"b_ptrs\"(#loc25))\n#loc76 = loc(callsite(#loc4 at #loc26))\n#loc77 = loc(callsite(#loc6 at #loc26))\n#loc78 = loc(\"a_ptrs\"(#loc2))\n#loc79 = loc(\"a_block\"(#loc27))\n#loc80 = loc(\"a_block\"(#loc28))\n#loc81 = loc(\"a_block\"(#loc29))\n#loc82 = loc(\"a_block\"(#loc30))\n#loc83 = loc(\"b_block\"(#loc31))\n#loc84 = loc(\"b_block\"(#loc32))\n#loc85 = loc(\"accumulator\"(#loc33))\n#loc86 = loc(\"a_ptrs\"(#loc34))\n#loc87 = loc(\"b_ptrs\"(#loc35))\n#loc88 = loc(\"b_ptrs\"(#loc36))\n#loc89 = loc(\"c_block\"(#loc37))\n#loc90 = loc(\"c_ptrs\"(#loc38))\n#loc91 = loc(\"c_ptrs\"(#loc39))\n#loc92 = loc(\"c_ptrs\"(#loc40))\n#loc93 = loc(\"c_ptrs\"(#loc41))\n#loc94 = loc(\"c_mask\"(#loc42))\n#loc95 = loc(\"c_mask\"(#loc43))\n#loc96 = loc(\"c_mask\"(#loc44))\n#loc97 = loc(callsite(#loc4 at #loc56))\n#loc98 = loc(callsite(#loc6 at #loc56))\n#loc99 = loc(callsite(#loc4 at #loc57))\n#loc100 = loc(callsite(#loc6 at #loc57))\n#loc101 = loc(\"b_ptrs\"(#loc78))\n#loc102 = loc(\"accumulator\"(#loc101))\n","matmul_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1, 8], threadsPerWarp = [16, 2], warpsPerCTA = [1, 1], order = [1, 0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:5)\n#mma = #ttg.nvidia_mma<{versionMajor = 2, versionMinor = 0, warpsPerCTA = [1, 1], instrShape = [16, 8]}>\n#shared = #ttg.swizzled_shared<{vec = 8, perPhase = 4, maxPhase = 2, order = [1, 0]}>\n#smem = #ttg.shared_memory\n#loc46 = loc(\"a\"(#loc))\n#loc47 = loc(\"b\"(#loc))\n#loc48 = loc(\"c\"(#loc))\n#loc49 = loc(\"M\"(#loc))\n#loc50 = loc(\"N\"(#loc))\n#loc51 = loc(\"K\"(#loc))\n#loc52 = loc(\"stride_am\"(#loc))\n#loc53 = loc(\"stride_bk\"(#loc))\n#loc54 = loc(\"stride_cm\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 1 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %cst = arith.constant dense<0.000000e+00> : tensor<16x16xf32, #mma> loc(#loc1)\n %cst_0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #blocked> loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc1)\n %cst_1 = arith.constant dense<16> : tensor<16x16xi32, #blocked> loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc55)\n %num_pid_m = arith.addi %M, %c15_i32 : i32 loc(#loc97)\n %num_pid_m_2 = arith.divsi %num_pid_m, %c16_i32 : i32 loc(#loc98)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc99)\n %num_pid_n_3 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc100)\n %group_id = arith.divsi %pid, %num_pid_n_3 : i32 loc(#loc58)\n %group_size = arith.subi %num_pid_m_2, %group_id : i32 loc(#loc59)\n %group_size_4 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc60)\n %pid_m = arith.remsi %pid, %group_size_4 : i32 loc(#loc61)\n %pid_m_5 = arith.addi %group_id, %pid_m : i32 loc(#loc62)\n %pid_n = arith.remsi %pid, %num_pid_n_3 : i32 loc(#loc63)\n %pid_n_6 = arith.divsi %pid_n, %group_size_4 : i32 loc(#loc64)\n %offs_am = arith.muli %pid_m_5, %c16_i32 : i32 loc(#loc65)\n %offs_am_7 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc66)\n %offs_am_8 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc66)\n %offs_am_9 = tt.splat %offs_am : i32 -> tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc65)\n %offs_am_10 = arith.addi %offs_am_9, %offs_am_7 : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc65)\n %offs_am_11 = tt.splat %M : i32 -> tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc67)\n %offs_am_12 = arith.remsi %offs_am_10, %offs_am_11 : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc67)\n %offs_bn = arith.muli %pid_n_6, %c16_i32 : i32 loc(#loc68)\n %offs_bn_13 = tt.splat %offs_bn : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc68)\n %offs_bn_14 = arith.addi %offs_bn_13, %offs_am_8 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc68)\n %offs_bn_15 = tt.splat %N : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc69)\n %offs_bn_16 = arith.remsi %offs_bn_14, %offs_bn_15 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc69)\n %a_ptrs = tt.expand_dims %offs_am_12 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc70)\n %a_ptrs_17 = tt.splat %stride_am : i32 -> tensor<16x1xi32, #blocked> loc(#loc70)\n %a_ptrs_18 = arith.muli %a_ptrs, %a_ptrs_17 : tensor<16x1xi32, #blocked> loc(#loc70)\n %a_ptrs_19 = tt.expand_dims %offs_am_8 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc71)\n %a_ptrs_20 = tt.broadcast %a_ptrs_18 : tensor<16x1xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc70)\n %a_ptrs_21 = tt.broadcast %a_ptrs_19 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc70)\n %a_ptrs_22 = arith.addi %a_ptrs_20, %a_ptrs_21 : tensor<16x16xi32, #blocked> loc(#loc70)\n %a_ptrs_23 = tt.splat %a : !tt.ptr -> tensor<16x16x!tt.ptr, #blocked> loc(#loc72)\n %a_ptrs_24 = tt.addptr %a_ptrs_23, %a_ptrs_22 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc72)\n %b_ptrs = tt.expand_dims %offs_am_7 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc73)\n %b_ptrs_25 = tt.splat %stride_bk : i32 -> tensor<16x1xi32, #blocked> loc(#loc73)\n %b_ptrs_26 = arith.muli %b_ptrs, %b_ptrs_25 : tensor<16x1xi32, #blocked> loc(#loc73)\n %b_ptrs_27 = tt.expand_dims %offs_bn_16 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc74)\n %b_ptrs_28 = tt.broadcast %b_ptrs_26 : tensor<16x1xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc73)\n %b_ptrs_29 = tt.broadcast %b_ptrs_27 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc73)\n %b_ptrs_30 = arith.addi %b_ptrs_28, %b_ptrs_29 : tensor<16x16xi32, #blocked> loc(#loc73)\n %b_ptrs_31 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr, #blocked> loc(#loc75)\n %b_ptrs_32 = tt.addptr %b_ptrs_31, %b_ptrs_30 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc75)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc76)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc77)\n %b_ptrs_33 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc78)\n %b_ptrs_34 = tt.splat %b_ptrs_33 : i32 -> tensor<16x16xi32, #blocked> loc(#loc79)\n %accumulator:3 = scf.for %k = %c0_i32 to %1 step %c1_i32 iter_args(%accumulator_49 = %cst, %a_ptrs_50 = %a_ptrs_24, %b_ptrs_51 = %b_ptrs_32) -> (tensor<16x16xf32, #mma>, tensor<16x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked>) : i32 {\n %a_block = arith.muli %k, %c16_i32 : i32 loc(#loc81)\n %a_block_52 = arith.subi %K, %a_block : i32 loc(#loc82)\n %a_block_53 = tt.splat %a_block_52 : i32 -> tensor<1x16xi32, #blocked> loc(#loc83)\n %a_block_54 = arith.cmpi slt, %a_ptrs_19, %a_block_53 : tensor<1x16xi32, #blocked> loc(#loc83)\n %a_block_55 = tt.broadcast %a_block_54 : tensor<1x16xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc84)\n %a_block_56 = tt.load %a_ptrs_50, %a_block_55, %cst_0 : tensor<16x16x!tt.ptr, #blocked> loc(#loc84)\n %a_block_57 = ttg.local_alloc %a_block_56 : (tensor<16x16xf16, #blocked>) -> !ttg.memdesc<16x16xf16, #shared, #smem> loc(#loc84)\n %b_block = tt.splat %a_block_52 : i32 -> tensor<16x1xi32, #blocked> loc(#loc85)\n %b_block_58 = arith.cmpi slt, %b_ptrs, %b_block : tensor<16x1xi32, #blocked> loc(#loc85)\n %b_block_59 = tt.broadcast %b_block_58 : tensor<16x1xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc86)\n %b_block_60 = tt.load %b_ptrs_51, %b_block_59, %cst_0 : tensor<16x16x!tt.ptr, #blocked> loc(#loc86)\n %b_block_61 = ttg.local_alloc %b_block_60 : (tensor<16x16xf16, #blocked>) -> !ttg.memdesc<16x16xf16, #shared, #smem> loc(#loc86)\n %a_block_62 = ttg.local_load %a_block_57 : !ttg.memdesc<16x16xf16, #shared, #smem> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> loc(#loc84)\n %b_block_63 = ttg.local_load %b_block_61 : !ttg.memdesc<16x16xf16, #shared, #smem> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> loc(#loc86)\n %accumulator_64 = tt.dot %a_block_62, %b_block_63, %accumulator_49, inputPrecision = tf32 : tensor<16x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> * tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> -> tensor<16x16xf32, #mma> loc(#loc87)\n %a_ptrs_65 = tt.addptr %a_ptrs_50, %cst_1 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc88)\n %b_ptrs_66 = tt.addptr %b_ptrs_51, %b_ptrs_34 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc79)\n scf.yield %accumulator_64, %a_ptrs_65, %b_ptrs_66 : tensor<16x16xf32, #mma>, tensor<16x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked> loc(#loc28)\n } loc(#loc102)\n %c_block = arith.truncf %accumulator#0 : tensor<16x16xf32, #mma> to tensor<16x16xf16, #mma> loc(#loc89)\n %c_ptrs = tt.expand_dims %offs_am_10 {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc90)\n %c_ptrs_35 = tt.splat %stride_cm : i32 -> tensor<16x1xi32, #blocked> loc(#loc91)\n %c_ptrs_36 = arith.muli %c_ptrs_35, %c_ptrs : tensor<16x1xi32, #blocked> loc(#loc91)\n %c_ptrs_37 = tt.splat %c : !tt.ptr -> tensor<16x1x!tt.ptr, #blocked> loc(#loc92)\n %c_ptrs_38 = tt.addptr %c_ptrs_37, %c_ptrs_36 : tensor<16x1x!tt.ptr, #blocked>, tensor<16x1xi32, #blocked> loc(#loc92)\n %c_ptrs_39 = tt.expand_dims %offs_bn_14 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc93)\n %c_ptrs_40 = tt.broadcast %c_ptrs_38 : tensor<16x1x!tt.ptr, #blocked> -> tensor<16x16x!tt.ptr, #blocked> loc(#loc92)\n %c_ptrs_41 = tt.broadcast %c_ptrs_39 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc92)\n %c_ptrs_42 = tt.addptr %c_ptrs_40, %c_ptrs_41 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc92)\n %c_mask = tt.splat %M : i32 -> tensor<16x1xi32, #blocked> loc(#loc94)\n %c_mask_43 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<16x1xi32, #blocked> loc(#loc94)\n %c_mask_44 = tt.splat %N : i32 -> tensor<1x16xi32, #blocked> loc(#loc95)\n %c_mask_45 = arith.cmpi slt, %c_ptrs_39, %c_mask_44 : tensor<1x16xi32, #blocked> loc(#loc95)\n %c_mask_46 = tt.broadcast %c_mask_43 : tensor<16x1xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc96)\n %c_mask_47 = tt.broadcast %c_mask_45 : tensor<1x16xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc96)\n %c_mask_48 = arith.andi %c_mask_46, %c_mask_47 : tensor<16x16xi1, #blocked> loc(#loc96)\n %2 = ttg.convert_layout %c_block : tensor<16x16xf16, #mma> -> tensor<16x16xf16, #blocked> loc(#loc45)\n tt.store %c_ptrs_42, %2, %c_mask_48 : tensor<16x16x!tt.ptr, #blocked> loc(#loc45)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:15)\n#loc3 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:13)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:21)\n#loc5 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:12)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:20)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:26)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:22)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:32)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:17)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:18)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:17)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:20)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:43)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:19)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:20)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:19)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:23)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:54)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:23)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:53)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:27)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:23)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:13)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:9)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:52)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:48)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:30)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":329:23)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:30)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":332:23)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:28)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:13)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:19)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:34)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:22)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:18)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:65)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:19)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:44)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:18)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:9)\n#loc55 = loc(\"pid\"(#loc2))\n#loc56 = loc(\"num_pid_m\"(#loc4))\n#loc57 = loc(\"num_pid_n\"(#loc6))\n#loc58 = loc(\"group_id\"(#loc7))\n#loc59 = loc(\"group_size\"(#loc8))\n#loc60 = loc(\"group_size\"(#loc9))\n#loc61 = loc(\"pid_m\"(#loc10))\n#loc62 = loc(\"pid_m\"(#loc11))\n#loc63 = loc(\"pid_n\"(#loc12))\n#loc64 = loc(\"pid_n\"(#loc13))\n#loc65 = loc(\"offs_am\"(#loc14))\n#loc66 = loc(\"offs_am\"(#loc15))\n#loc67 = loc(\"offs_am\"(#loc16))\n#loc68 = loc(\"offs_bn\"(#loc17))\n#loc69 = loc(\"offs_bn\"(#loc18))\n#loc70 = loc(\"a_ptrs\"(#loc19))\n#loc71 = loc(\"a_ptrs\"(#loc20))\n#loc72 = loc(\"a_ptrs\"(#loc21))\n#loc73 = loc(\"b_ptrs\"(#loc22))\n#loc74 = loc(\"b_ptrs\"(#loc23))\n#loc75 = loc(\"b_ptrs\"(#loc24))\n#loc76 = loc(callsite(#loc3 at #loc25))\n#loc77 = loc(callsite(#loc5 at #loc25))\n#loc78 = loc(\"b_ptrs\"(#loc26))\n#loc79 = loc(\"b_ptrs\"(#loc27))\n#loc80 = loc(\"a_ptrs\"(#loc28))\n#loc81 = loc(\"a_block\"(#loc29))\n#loc82 = loc(\"a_block\"(#loc30))\n#loc83 = loc(\"a_block\"(#loc31))\n#loc84 = loc(\"a_block\"(#loc32))\n#loc85 = loc(\"b_block\"(#loc33))\n#loc86 = loc(\"b_block\"(#loc34))\n#loc87 = loc(\"accumulator\"(#loc35))\n#loc88 = loc(\"a_ptrs\"(#loc36))\n#loc89 = loc(\"c_block\"(#loc37))\n#loc90 = loc(\"c_ptrs\"(#loc38))\n#loc91 = loc(\"c_ptrs\"(#loc39))\n#loc92 = loc(\"c_ptrs\"(#loc40))\n#loc93 = loc(\"c_ptrs\"(#loc41))\n#loc94 = loc(\"c_mask\"(#loc42))\n#loc95 = loc(\"c_mask\"(#loc43))\n#loc96 = loc(\"c_mask\"(#loc44))\n#loc97 = loc(callsite(#loc3 at #loc56))\n#loc98 = loc(callsite(#loc5 at #loc56))\n#loc99 = loc(callsite(#loc3 at #loc57))\n#loc100 = loc(callsite(#loc5 at #loc57))\n#loc101 = loc(\"b_ptrs\"(#loc80))\n#loc102 = loc(\"accumulator\"(#loc101))\n","matmul_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n@global_smem = external addrspace(3) global [0 x i8], align 16\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @matmul_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i32 %8, ptr addrspace(1) readnone captures(none) %9, ptr addrspace(1) readnone captures(none) %10) local_unnamed_addr #0 !dbg !4 {\n %12 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %13 = add i32 %3, 15, !dbg !13\n %14 = sdiv i32 %13, 16, !dbg !18\n %15 = add i32 %4, 15, !dbg !19\n %16 = sdiv i32 %15, 16, !dbg !21\n %.frozen = freeze i32 %16, !dbg !22\n %17 = sdiv i32 %12, %.frozen, !dbg !22\n %18 = sub i32 %14, %17, !dbg !23\n %19 = tail call i32 @llvm.smin.i32(i32 %18, i32 1), !dbg !24\n %20 = srem i32 %12, %19, !dbg !25\n %21 = add i32 %20, %17, !dbg !26\n %22 = mul i32 %17, %.frozen, !dbg !27\n %.decomposed = sub i32 %12, %22, !dbg !27\n %23 = sdiv i32 %.decomposed, %19, !dbg !28\n %24 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !29\n %25 = lshr i32 %24, 1, !dbg !29\n %26 = and i32 %24, 1, !dbg !29\n %27 = and i32 %25, 15, !dbg !29\n %28 = shl nuw nsw i32 %26, 3, !dbg !29\n %29 = insertelement <2 x i32> poison, i32 %23, i64 0, !dbg !30\n %30 = insertelement <2 x i32> %29, i32 %21, i64 1, !dbg !30\n %31 = shl <2 x i32> %30, splat (i32 4), !dbg !30\n %32 = insertelement <2 x i32> poison, i32 %28, i64 0, !dbg !30\n %33 = insertelement <2 x i32> %32, i32 %27, i64 1, !dbg !30\n %34 = or disjoint <2 x i32> %31, %33, !dbg !30\n %35 = add i32 %5, 15, !dbg !31\n %36 = sdiv i32 %35, 16, !dbg !33\n %37 = icmp sgt i32 %35, 15, !dbg !34\n br i1 %37, label %.lr.ph, label %.._crit_edge_crit_edge, !dbg !34\n\n.._crit_edge_crit_edge: ; preds = %11\n %.pre = shl nuw nsw i32 %24, 4, !dbg !35\n %.pre65 = and i32 %24, 4, !dbg !35\n %.pre67 = and i32 %24, 8, !dbg !35\n br label %._crit_edge, !dbg !34\n\n.lr.ph: ; preds = %11\n %38 = shl i32 %7, 4, !dbg !36\n %39 = mul i32 %7, %27, !dbg !37\n %40 = insertelement <2 x i32> poison, i32 %4, i64 0, !dbg !38\n %41 = insertelement <2 x i32> %40, i32 %3, i64 1, !dbg !38\n %42 = srem <2 x i32> %34, %41, !dbg !38\n %43 = extractelement <2 x i32> %42, i64 0, !dbg !37\n %44 = add i32 %43, %39, !dbg !37\n %45 = sext i32 %44 to i64, !dbg !39\n %46 = getelementptr [2 x i8], ptr addrspace(1) %1, i64 %45, !dbg !39\n %47 = extractelement <2 x i32> %42, i64 1, !dbg !40\n %48 = mul i32 %47, %6, !dbg !40\n %49 = add i32 %48, %28, !dbg !40\n %50 = sext i32 %49 to i64, !dbg !41\n %51 = getelementptr [2 x i8], ptr addrspace(1) %0, i64 %50, !dbg !41\n %52 = shl nuw nsw i32 %24, 4\n %53 = and i32 %52, 368\n %54 = and i32 %24, 8\n %55 = icmp eq i32 %54, 0\n %56 = select i1 %55, i32 0, i32 144\n %57 = xor i32 %56, %53\n %58 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %57\n %59 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 512), i32 %57\n %60 = shl nuw nsw i32 %24, 5\n %61 = and i32 %60, 352\n %62 = and i32 %24, 4\n %63 = icmp eq i32 %62, 0\n %64 = select i1 %63, i32 0, i32 144\n %65 = and i32 %24, 16\n %66 = xor i32 %64, %65\n %67 = or disjoint i32 %66, %61\n %68 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %67\n %69 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 512), i32 %67\n %70 = sext i32 %38 to i64\n br label %71, !dbg !34\n\n71: ; preds = %.lr.ph, %71\n %.pn3163 = phi ptr addrspace(1) [ %46, %.lr.ph ], [ %116, %71 ]\n %.pn1562 = phi ptr addrspace(1) [ %51, %.lr.ph ], [ %115, %71 ]\n %72 = phi float [ 0.000000e+00, %.lr.ph ], [ %106, %71 ]\n %73 = phi float [ 0.000000e+00, %.lr.ph ], [ %107, %71 ]\n %74 = phi float [ 0.000000e+00, %.lr.ph ], [ %108, %71 ]\n %75 = phi float [ 0.000000e+00, %.lr.ph ], [ %109, %71 ]\n %76 = phi float [ 0.000000e+00, %.lr.ph ], [ %111, %71 ]\n %77 = phi float [ 0.000000e+00, %.lr.ph ], [ %112, %71 ]\n %78 = phi float [ 0.000000e+00, %.lr.ph ], [ %113, %71 ]\n %79 = phi float [ 0.000000e+00, %.lr.ph ], [ %114, %71 ]\n %80 = phi i32 [ 0, %.lr.ph ], [ %117, %71 ]\n %81 = shl i32 %80, 4, !dbg !42\n %82 = sub i32 %5, %81, !dbg !43\n %83 = icmp slt i32 %28, %82, !dbg !44\n %84 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn1562, i1 %83) #6, !dbg !45\n %85 = extractvalue { i32, i32, i32, i32 } %84, 0, !dbg !45\n %86 = extractvalue { i32, i32, i32, i32 } %84, 1, !dbg !45\n %87 = extractvalue { i32, i32, i32, i32 } %84, 2, !dbg !45\n %88 = extractvalue { i32, i32, i32, i32 } %84, 3, !dbg !45\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !45\n tail call void asm sideeffect \"st.shared::cta.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %58, i32 %85, i32 %86, i32 %87, i32 %88) #6, !dbg !45\n %89 = icmp slt i32 %27, %82, !dbg !46\n %90 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn3163, i1 %89) #6, !dbg !47\n %91 = extractvalue { i32, i32, i32, i32 } %90, 0, !dbg !47\n %92 = extractvalue { i32, i32, i32, i32 } %90, 1, !dbg !47\n %93 = extractvalue { i32, i32, i32, i32 } %90, 2, !dbg !47\n %94 = extractvalue { i32, i32, i32, i32 } %90, 3, !dbg !47\n tail call void asm sideeffect \"st.shared::cta.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %59, i32 %91, i32 %92, i32 %93, i32 %94) #6, !dbg !47\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !45\n %95 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %68), !dbg !45\n %96 = extractvalue { i32, i32, i32, i32 } %95, 0, !dbg !45\n %97 = extractvalue { i32, i32, i32, i32 } %95, 1, !dbg !45\n %98 = extractvalue { i32, i32, i32, i32 } %95, 2, !dbg !45\n %99 = extractvalue { i32, i32, i32, i32 } %95, 3, !dbg !45\n %100 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %69), !dbg !47\n %101 = extractvalue { i32, i32, i32, i32 } %100, 0, !dbg !47\n %102 = extractvalue { i32, i32, i32, i32 } %100, 1, !dbg !47\n %103 = extractvalue { i32, i32, i32, i32 } %100, 2, !dbg !47\n %104 = extractvalue { i32, i32, i32, i32 } %100, 3, !dbg !47\n %105 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %72, float %73, float %74, float %75, i32 %96, i32 %97, i32 %98, i32 %99, i32 %101, i32 %102) #6, !dbg !48\n %106 = extractvalue { float, float, float, float } %105, 0, !dbg !48\n %107 = extractvalue { float, float, float, float } %105, 1, !dbg !48\n %108 = extractvalue { float, float, float, float } %105, 2, !dbg !48\n %109 = extractvalue { float, float, float, float } %105, 3, !dbg !48\n %110 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %76, float %77, float %78, float %79, i32 %96, i32 %97, i32 %98, i32 %99, i32 %103, i32 %104) #6, !dbg !48\n %111 = extractvalue { float, float, float, float } %110, 0, !dbg !48\n %112 = extractvalue { float, float, float, float } %110, 1, !dbg !48\n %113 = extractvalue { float, float, float, float } %110, 2, !dbg !48\n %114 = extractvalue { float, float, float, float } %110, 3, !dbg !48\n %115 = getelementptr i8, ptr addrspace(1) %.pn1562, i64 32, !dbg !49\n %116 = getelementptr [2 x i8], ptr addrspace(1) %.pn3163, i64 %70, !dbg !50\n %117 = add nuw nsw i32 %80, 1, !dbg !34\n %exitcond.not = icmp eq i32 %117, %36, !dbg !34\n br i1 %exitcond.not, label %._crit_edge.loopexit, label %71, !dbg !34\n\n._crit_edge.loopexit: ; preds = %71\n %118 = insertelement <8 x float> poison, float %106, i64 0, !dbg !51\n %119 = insertelement <8 x float> %118, float %107, i64 1, !dbg !51\n %120 = insertelement <8 x float> %119, float %108, i64 2, !dbg !51\n %121 = insertelement <8 x float> %120, float %109, i64 3, !dbg !51\n %122 = insertelement <8 x float> %121, float %111, i64 4, !dbg !51\n %123 = insertelement <8 x float> %122, float %112, i64 5, !dbg !51\n %124 = insertelement <8 x float> %123, float %113, i64 6, !dbg !51\n %125 = insertelement <8 x float> %124, float %114, i64 7, !dbg !51\n %126 = fptrunc <8 x float> %125 to <8 x half>, !dbg !51\n %127 = bitcast <8 x half> %126 to <4 x i32>, !dbg !35\n br label %._crit_edge, !dbg !51\n\n._crit_edge: ; preds = %._crit_edge.loopexit, %.._crit_edge_crit_edge\n %.pre-phi68 = phi i32 [ %.pre67, %.._crit_edge_crit_edge ], [ %54, %._crit_edge.loopexit ], !dbg !35\n %.pre-phi66 = phi i32 [ %.pre65, %.._crit_edge_crit_edge ], [ %62, %._crit_edge.loopexit ], !dbg !35\n %.pre-phi = phi i32 [ %.pre, %.._crit_edge_crit_edge ], [ %52, %._crit_edge.loopexit ], !dbg !35\n %.bc82 = phi <4 x i32> [ zeroinitializer, %.._crit_edge_crit_edge ], [ %127, %._crit_edge.loopexit ], !dbg !52\n %128 = extractelement <2 x i32> %34, i64 1, !dbg !53\n %129 = mul i32 %128, %8, !dbg !53\n %130 = sext i32 %129 to i64, !dbg !54\n %131 = getelementptr [2 x i8], ptr addrspace(1) %2, i64 %130, !dbg !54\n %132 = extractelement <2 x i32> %34, i64 0, !dbg !54\n %133 = sext i32 %132 to i64, !dbg !54\n %134 = getelementptr [2 x i8], ptr addrspace(1) %131, i64 %133, !dbg !54\n %135 = icmp slt i32 %128, %3, !dbg !55\n %136 = icmp slt i32 %132, %4, !dbg !56\n %137 = and i1 %135, %136, !dbg !57\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !35\n %138 = and i32 %.pre-phi, 176, !dbg !35\n %139 = icmp eq i32 %.pre-phi66, 0, !dbg !35\n %140 = select i1 %139, i32 0, i32 320, !dbg !35\n %141 = shl nuw nsw i32 %24, 2, !dbg !35\n %142 = and i32 %141, 64, !dbg !35\n %143 = or disjoint i32 %140, %138, !dbg !35\n %144 = xor i32 %143, %142, !dbg !35\n %145 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %144, !dbg !35\n %.extract77 = extractelement <4 x i32> %.bc82, i64 0, !dbg !35\n %.extract79 = extractelement <4 x i32> %.bc82, i64 1, !dbg !35\n %.extract81 = extractelement <4 x i32> %.bc82, i64 2, !dbg !35\n %.extract83 = extractelement <4 x i32> %.bc82, i64 3, !dbg !35\n tail call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %145, i32 %.extract77, i32 %.extract79, i32 %.extract81, i32 %.extract83), !dbg !35\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !35\n %146 = shl nuw nsw i32 %24, 3, !dbg !35\n %147 = and i32 %146, 176, !dbg !35\n %148 = shl nuw nsw i32 %26, 6, !dbg !35\n %149 = icmp eq i32 %.pre-phi68, 0, !dbg !35\n %150 = select i1 %149, i32 0, i32 320, !dbg !35\n %151 = or disjoint i32 %147, %148, !dbg !35\n %152 = xor i32 %151, %150, !dbg !35\n %153 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %152, !dbg !35\n %.extract = load i32, ptr addrspace(3) %153, align 16, !dbg !35\n %154 = getelementptr inbounds nuw i8, ptr addrspace(3) %153, i32 4, !dbg !35\n %.extract57 = load i32, ptr addrspace(3) %154, align 4, !dbg !35\n %155 = getelementptr inbounds nuw i8, ptr addrspace(3) %153, i32 8, !dbg !35\n %.extract59 = load i32, ptr addrspace(3) %155, align 8, !dbg !35\n %156 = getelementptr inbounds nuw i8, ptr addrspace(3) %153, i32 12, !dbg !35\n %.extract61 = load i32, ptr addrspace(3) %156, align 4, !dbg !35\n tail call void asm sideeffect \"@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };\", \"r,r,r,r,l,b\"(i32 %.extract, i32 %.extract57, i32 %.extract59, i32 %.extract61, ptr addrspace(1) %134, i1 %137) #6, !dbg !35\n ret void, !dbg !58\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)\ndeclare i32 @llvm.smin.i32(i32, i32) #2\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3\n\n; Function Attrs: convergent nocallback nounwind memory(argmem: write)\ndeclare void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) writeonly captures(none), i32, i32, i32, i32) #4\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.bar.warp.sync(i32) #3\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"32\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #3 = { convergent nocallback nounwind }\nattributes #4 = { convergent nocallback nounwind memory(argmem: write) }\nattributes #5 = { convergent nocallback nofree nounwind memory(argmem: read) }\nattributes #6 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"matmul_kernel\", linkageName: \"matmul_kernel\", scope: !1, file: !1, line: 292, type: !5, scopeLine: 292, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !9, !9, !9, !9, !9, !9, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"half\", size: 16, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 311, column: 15, scope: !4)\n!13 = !DILocation(line: 43, column: 13, scope: !14, inlinedAt: !16)\n!14 = distinct !DILexicalBlockFile(scope: !4, file: !15, discriminator: 0)\n!15 = !DIFile(filename: \"standard.py\", directory: \"/python/lib/python3.14/site-packages/triton/language\")\n!16 = !DILocation(line: 312, column: 21, scope: !17)\n!17 = distinct !DILexicalBlockFile(scope: !4, file: !1, discriminator: 0)\n!18 = !DILocation(line: 43, column: 12, scope: !14, inlinedAt: !16)\n!19 = !DILocation(line: 43, column: 13, scope: !14, inlinedAt: !20)\n!20 = !DILocation(line: 313, column: 21, scope: !17)\n!21 = !DILocation(line: 43, column: 12, scope: !14, inlinedAt: !20)\n!22 = !DILocation(line: 315, column: 20, scope: !4)\n!23 = !DILocation(line: 317, column: 26, scope: !4)\n!24 = !DILocation(line: 317, column: 22, scope: !4)\n!25 = !DILocation(line: 318, column: 32, scope: !4)\n!26 = !DILocation(line: 318, column: 17, scope: !4)\n!27 = !DILocation(line: 319, column: 18, scope: !4)\n!28 = !DILocation(line: 319, column: 17, scope: !4)\n!29 = !DILocation(line: 321, column: 43, scope: !4)\n!30 = !DILocation(line: 322, column: 20, scope: !4)\n!31 = !DILocation(line: 43, column: 13, scope: !14, inlinedAt: !32)\n!32 = !DILocation(line: 328, column: 27, scope: !17)\n!33 = !DILocation(line: 43, column: 12, scope: !14, inlinedAt: !32)\n!34 = !DILocation(line: 328, column: 9, scope: !4)\n!35 = !DILocation(line: 344, column: 9, scope: !4)\n!36 = !DILocation(line: 337, column: 23, scope: !4)\n!37 = !DILocation(line: 325, column: 23, scope: !4)\n!38 = !DILocation(line: 322, column: 19, scope: !4)\n!39 = !DILocation(line: 325, column: 18, scope: !4)\n!40 = !DILocation(line: 324, column: 23, scope: !4)\n!41 = !DILocation(line: 324, column: 18, scope: !4)\n!42 = !DILocation(line: 330, column: 52, scope: !4)\n!43 = !DILocation(line: 330, column: 48, scope: !4)\n!44 = !DILocation(line: 330, column: 30, scope: !4)\n!45 = !DILocation(line: 329, column: 23, scope: !4)\n!46 = !DILocation(line: 333, column: 30, scope: !4)\n!47 = !DILocation(line: 332, column: 23, scope: !4)\n!48 = !DILocation(line: 335, column: 28, scope: !4)\n!49 = !DILocation(line: 336, column: 13, scope: !4)\n!50 = !DILocation(line: 337, column: 13, scope: !4)\n!51 = !DILocation(line: 338, column: 19, scope: !4)\n!52 = !DILocation(line: 327, column: 23, scope: !4)\n!53 = !DILocation(line: 342, column: 22, scope: !4)\n!54 = !DILocation(line: 342, column: 18, scope: !4)\n!55 = !DILocation(line: 343, column: 19, scope: !4)\n!56 = !DILocation(line: 343, column: 44, scope: !4)\n!57 = !DILocation(line: 343, column: 18, scope: !4)\n!58 = !DILocation(line: 292, column: 5, scope: !4)\n","matmul_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.8\n.target sm_90a\n.address_size 64\n\n\t// .globl\tmatmul_kernel // -- Begin function matmul_kernel\n.extern .shared .align 16 .b8 global_smem[];\n // @matmul_kernel\n.visible .entry matmul_kernel(\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_2,\n\t.param .u32 matmul_kernel_param_3,\n\t.param .u32 matmul_kernel_param_4,\n\t.param .u32 matmul_kernel_param_5,\n\t.param .u32 matmul_kernel_param_6,\n\t.param .u32 matmul_kernel_param_7,\n\t.param .u32 matmul_kernel_param_8,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_9,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_10\n)\n.reqntid 32\n{\n\t.reg .pred \t%p<10>;\n\t.reg .b32 \t%r<111>;\n\t.reg .b64 \t%rd<9>;\n\t.loc\t1 292 0 // generate_examples.py:292:0\n$L__func_begin0:\n\t.loc\t1 292 0 // generate_examples.py:292:0\n\n// %bb.0:\n\tld.param.b32 \t%r13, [matmul_kernel_param_8];\n\tld.param.b32 \t%r94, [matmul_kernel_param_5];\n\tld.param.b32 \t%r10, [matmul_kernel_param_4];\n\tld.param.b32 \t%r9, [matmul_kernel_param_3];\n\tld.param.b64 \t%rd4, [matmul_kernel_param_2];\n$L__tmp0:\n\t.loc\t1 311 15 // generate_examples.py:311:15\n\tmov.u32 \t%r14, %ctaid.x;\n$L__tmp1:\n\t.loc\t1 312 21 // generate_examples.py:312:21\n\t.loc\t2 43 13, function_name $L__info_string0, inlined_at 1 312 21 // standard.py:43:13\n\tadd.s32 \t%r15, %r9, 15;\n\t.loc\t2 43 12, function_name $L__info_string0, inlined_at 1 312 21 // standard.py:43:12\n\tshr.s32 \t%r16, %r15, 31;\n\tshr.u32 \t%r17, %r16, 28;\n\tadd.s32 \t%r18, %r15, %r17;\n\tshr.s32 \t%r19, %r18, 4;\n$L__tmp2:\n\t.loc\t1 313 21 // generate_examples.py:313:21\n\t.loc\t2 43 13, function_name $L__info_string0, inlined_at 1 313 21 // standard.py:43:13\n\tadd.s32 \t%r20, %r10, 15;\n\t.loc\t2 43 12, function_name $L__info_string0, inlined_at 1 313 21 // standard.py:43:12\n\tshr.s32 \t%r21, %r20, 31;\n\tshr.u32 \t%r22, %r21, 28;\n\tadd.s32 \t%r23, %r20, %r22;\n\tshr.s32 \t%r24, %r23, 4;\n$L__tmp3:\n\t.loc\t1 315 20 // generate_examples.py:315:20\n\tdiv.s32 \t%r25, %r14, %r24;\n\t.loc\t1 317 26 // generate_examples.py:317:26\n\tsub.s32 \t%r26, %r19, %r25;\n\t.loc\t1 317 22 // generate_examples.py:317:22\n\tmin.s32 \t%r27, %r26, 1;\n\t.loc\t1 318 32 // generate_examples.py:318:32\n\trem.s32 \t%r28, %r14, %r27;\n\t.loc\t1 318 17 // generate_examples.py:318:17\n\tadd.s32 \t%r29, %r28, %r25;\n\t.loc\t1 319 18 // generate_examples.py:319:18\n\tmul.lo.s32 \t%r30, %r25, %r24;\n\tsub.s32 \t%r31, %r14, %r30;\n\t.loc\t1 319 17 // generate_examples.py:319:17\n\tdiv.s32 \t%r32, %r31, %r27;\n\t.loc\t1 321 43 // generate_examples.py:321:43\n\tmov.u32 \t%r1, %tid.x;\n\tand.b32 \t%r2, %r1, 1;\n\tbfe.u32 \t%r3, %r1, 1, 4;\n\tshl.b32 \t%r4, %r2, 3;\n\t.loc\t1 322 20 // generate_examples.py:322:20\n\tshl.b32 \t%r33, %r32, 4;\n\tshl.b32 \t%r34, %r29, 4;\n\tor.b32 \t%r5, %r33, %r4;\n\tor.b32 \t%r6, %r34, %r3;\n$L__tmp4:\n\t.loc\t1 328 27 // generate_examples.py:328:27\n\t.loc\t2 43 13, function_name $L__info_string0, inlined_at 1 328 27 // standard.py:43:13\n\tadd.s32 \t%r35, %r94, 15;\n$L__tmp5:\n\t.loc\t1 328 9 // generate_examples.py:328:9\n\tsetp.gt.s32 \t%p1, %r35, 15;\n\tmov.b32 \t%r93, global_smem;\n\t@!%p1 bra \t$L__BB0_1;\n// %bb.2: // %.lr.ph\n\t.loc\t1 0 9 // generate_examples.py:0:9\n\tld.param.b32 \t%r12, [matmul_kernel_param_7];\n\tld.param.b32 \t%r11, [matmul_kernel_param_6];\n\tld.param.b64 \t%rd3, [matmul_kernel_param_1];\n\tld.param.b64 \t%rd2, [matmul_kernel_param_0];\n\tshr.s32 \t%r36, %r35, 31;\n\tshr.u32 \t%r37, %r36, 28;\n\tadd.s32 \t%r38, %r35, %r37;\n\tshr.s32 \t%r95, %r38, 4;\n\t.loc\t1 337 23 // generate_examples.py:337:23\n\tshl.b32 \t%r39, %r12, 4;\n\t.loc\t1 322 19 // generate_examples.py:322:19\n\trem.s32 \t%r40, %r6, %r9;\n\trem.s32 \t%r41, %r5, %r10;\n\t.loc\t1 325 23 // generate_examples.py:325:23\n\tmad.lo.s32 \t%r42, %r12, %r3, %r41;\n\t.loc\t1 325 18 // generate_examples.py:325:18\n\tmad.wide.s32 \t%rd7, %r42, 2, %rd3;\n\t.loc\t1 324 23 // generate_examples.py:324:23\n\tmad.lo.s32 \t%r43, %r40, %r11, %r4;\n\t.loc\t1 324 18 // generate_examples.py:324:18\n\tmad.wide.s32 \t%rd8, %r43, 2, %rd2;\n\tshl.b32 \t%r106, %r1, 4;\n\tand.b32 \t%r44, %r106, 368;\n\tand.b32 \t%r104, %r1, 8;\n\tbfe.s32 \t%r45, %r1, 3, 1;\n\tand.b32 \t%r46, %r45, 144;\n\txor.b32 \t%r47, %r46, %r44;\n\tadd.s32 \t%r60, %r93, %r47;\n\tadd.s32 \t%r65, %r60, 512;\n\tshl.b32 \t%r48, %r1, 5;\n\tand.b32 \t%r49, %r48, 352;\n\tand.b32 \t%r105, %r1, 4;\n\tbfe.s32 \t%r50, %r1, 2, 1;\n\tand.b32 \t%r51, %r50, 144;\n\tand.b32 \t%r52, %r1, 16;\n\txor.b32 \t%r53, %r51, %r52;\n\tor.b32 \t%r54, %r53, %r49;\n\tadd.s32 \t%r7, %r93, %r54;\n\tadd.s32 \t%r8, %r7, 512;\n\t.loc\t1 328 9 // generate_examples.py:328:9\n\tmul.wide.s32 \t%rd1, %r39, 2;\n\tmov.b32 \t%r96, 0f00000000;\n\tmov.b32 \t%r97, %r96;\n\tmov.b32 \t%r98, %r96;\n\tmov.b32 \t%r99, %r96;\n\tmov.b32 \t%r100, %r96;\n\tmov.b32 \t%r101, %r96;\n\tmov.b32 \t%r102, %r96;\n\tmov.b32 \t%r103, %r96;\n$L__BB0_3: // =>This Inner Loop Header: Depth=1\n\t.loc\t1 330 30 // generate_examples.py:330:30\n\tsetp.lt.s32 \t%p2, %r4, %r94;\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tmov.b32 \t%r59, 0;\n\t// begin inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tmov.u32 %r55, %r59;\n\t.loc\t1 330 23 // generate_examples.py:330:23\n\tmov.u32 %r56, %r59;\n\t.loc\t1 331 23 // generate_examples.py:331:23\n\tmov.u32 %r57, %r59;\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tmov.u32 %r58, %r59;\n\t@%p2 ld.global.v4.b32 { %r55, %r56, %r57, %r58 }, [ %rd8 + 0 ];\n\n\t// end inline asm\n\tbar.sync \t0;\n\t// begin inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tst.shared::cta.v4.b32 [ %r60 + 0 ], { %r55, %r56, %r57, %r58 };\n\n\t// end inline asm\n\t.loc\t1 333 30 // generate_examples.py:333:30\n\tsetp.lt.s32 \t%p3, %r3, %r94;\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\t// begin inline asm\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tmov.u32 %r61, %r59;\n\t.loc\t1 333 23 // generate_examples.py:333:23\n\tmov.u32 %r62, %r59;\n\t.loc\t1 334 23 // generate_examples.py:334:23\n\tmov.u32 %r63, %r59;\n\t.loc\t1 335 23 // generate_examples.py:335:23\n\tmov.u32 %r64, %r59;\n\t@%p3 ld.global.v4.b32 { %r61, %r62, %r63, %r64 }, [ %rd7 + 0 ];\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tst.shared::cta.v4.b32 [ %r65 + 0 ], { %r61, %r62, %r63, %r64 };\n\n\t// end inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tbar.sync \t0;\n\tldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r66, %r67, %r68, %r69}, [%r7];\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r70, %r71, %r72, %r73}, [%r8];\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\t// begin inline asm\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r96, %r97, %r98, %r99 }, { %r66, %r67, %r68, %r69 }, { %r70, %r71 }, { %r96, %r97, %r98, %r99 };\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r100, %r101, %r102, %r103 }, { %r66, %r67, %r68, %r69 }, { %r72, %r73 }, { %r100, %r101, %r102, %r103 };\n\n\t// end inline asm\n\t.loc\t1 336 13 // generate_examples.py:336:13\n\tadd.s64 \t%rd8, %rd8, 32;\n\t.loc\t1 328 9 // generate_examples.py:328:9\n\tadd.s64 \t%rd7, %rd7, %rd1;\n\tadd.s32 \t%r95, %r95, -1;\n\tadd.s32 \t%r94, %r94, -16;\n\tsetp.ne.b32 \t%p4, %r95, 0;\n\t@%p4 bra \t$L__BB0_3;\n// %bb.4: // %._crit_edge.loopexit\n\t.loc\t1 338 19 // generate_examples.py:338:19\n\tcvt.rn.f16x2.f32 \t%r110, %r103, %r102;\n\tcvt.rn.f16x2.f32 \t%r109, %r101, %r100;\n\tcvt.rn.f16x2.f32 \t%r108, %r99, %r98;\n\tcvt.rn.f16x2.f32 \t%r107, %r97, %r96;\n\tbra.uni \t$L__BB0_5;\n$L__BB0_1: // %.._crit_edge_crit_edge\n\t.loc\t1 344 9 // generate_examples.py:344:9\n\tshl.b32 \t%r106, %r1, 4;\n\tand.b32 \t%r105, %r1, 4;\n\tand.b32 \t%r104, %r1, 8;\n\tmov.b32 \t%r107, 0;\n\tmov.b32 \t%r108, %r107;\n\tmov.b32 \t%r109, %r107;\n\tmov.b32 \t%r110, %r107;\n$L__BB0_5: // %._crit_edge\n\t.loc\t1 342 22 // generate_examples.py:342:22\n\tmul.lo.s32 \t%r78, %r6, %r13;\n\t.loc\t1 342 18 // generate_examples.py:342:18\n\tmad.wide.s32 \t%rd6, %r78, 2, %rd4;\n\tmad.wide.s32 \t%rd5, %r5, 2, %rd6;\n\t.loc\t1 343 19 // generate_examples.py:343:19\n\tsetp.lt.s32 \t%p6, %r6, %r9;\n\t.loc\t1 343 44 // generate_examples.py:343:44\n\tsetp.lt.s32 \t%p7, %r5, %r10;\n\t.loc\t1 343 18 // generate_examples.py:343:18\n\tand.pred \t%p5, %p6, %p7;\n\t.loc\t1 344 9 // generate_examples.py:344:9\n\tbar.sync \t0;\n\tand.b32 \t%r79, %r106, 176;\n\tsetp.eq.b32 \t%p8, %r105, 0;\n\tselp.b32 \t%r80, 0, 320, %p8;\n\tshl.b32 \t%r81, %r1, 2;\n\tand.b32 \t%r82, %r81, 64;\n\tor.b32 \t%r83, %r80, %r79;\n\txor.b32 \t%r84, %r83, %r82;\n\tadd.s32 \t%r85, %r93, %r84;\n\tstmatrix.sync.aligned.m8n8.x4.shared.b16 [%r85], {%r107, %r108, %r109, %r110};\n\tbar.warp.sync \t-1;\n\tshl.b32 \t%r86, %r1, 3;\n\tand.b32 \t%r87, %r86, 176;\n\tshl.b32 \t%r88, %r2, 6;\n\tsetp.eq.b32 \t%p9, %r104, 0;\n\tselp.b32 \t%r89, 0, 320, %p9;\n\tor.b32 \t%r90, %r87, %r88;\n\txor.b32 \t%r91, %r90, %r89;\n\tadd.s32 \t%r92, %r93, %r91;\n\tld.shared.v4.b32 \t{%r74, %r75, %r76, %r77}, [%r92];\n\t// begin inline asm\n\t@%p5 st.global.v4.b32 [ %rd5 + 0 ], { %r74, %r75, %r76, %r77 };\n\n\t// end inline asm\n\t.loc\t1 292 5 // generate_examples.py:292:5\n\tret;\n$L__tmp6:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.file\t2 \"/python/lib/python3.14/site-packages/triton/language/standard.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 1 // DW_CHILDREN_yes\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 2 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 0 // DW_CHILDREN_no\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 32 // DW_AT_inline\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 3 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 1 // DW_CHILDREN_yes\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 4 // Abbreviation Code\n.b8 29 // DW_TAG_inlined_subroutine\n.b8 0 // DW_CHILDREN_no\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 88 // DW_AT_call_file\n.b8 11 // DW_FORM_data1\n.b8 89 // DW_AT_call_line\n.b8 5 // DW_FORM_data2\n.b8 87 // DW_AT_call_column\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 227 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0xdc DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n.b8 2 // Abbrev [2] 0x75:0x10 DW_TAG_subprogram\n.b8 109 // DW_AT_name\n.b8 97\n.b8 116\n.b8 109\n.b8 117\n.b8 108\n.b8 95\n.b8 107\n.b8 101\n.b8 114\n.b8 110\n.b8 101\n.b8 108\n.b8 0\n.b8 1 // DW_AT_inline\n.b8 3 // Abbrev [3] 0x85:0x61 DW_TAG_subprogram\n.b64 $L__func_begin0 // DW_AT_low_pc\n.b64 $L__func_end0 // DW_AT_high_pc\n.b32 117 // DW_AT_abstract_origin\n.b8 4 // Abbrev [4] 0x9a:0x19 DW_TAG_inlined_subroutine\n.b32 117 // DW_AT_abstract_origin\n.b64 $L__tmp1 // DW_AT_low_pc\n.b64 $L__tmp2 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 56 // DW_AT_call_line\n.b8 1\n.b8 21 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xb3:0x19 DW_TAG_inlined_subroutine\n.b32 117 // DW_AT_abstract_origin\n.b64 $L__tmp2 // DW_AT_low_pc\n.b64 $L__tmp3 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 57 // DW_AT_call_line\n.b8 1\n.b8 21 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xcc:0x19 DW_TAG_inlined_subroutine\n.b32 117 // DW_AT_abstract_origin\n.b64 $L__tmp4 // DW_AT_low_pc\n.b64 $L__tmp5 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 72 // DW_AT_call_line\n.b8 1\n.b8 27 // DW_AT_call_column\n.b8 0 // End Of Children Mark\n.b8 0 // End Of Children Mark\n\t}\n\t.section\t.debug_str\n\t{\n$L__info_string0:\n.b8 109 // string offset=0 ; matmul_kernel\n.b8 97\n.b8 116\n.b8 109\n.b8 117\n.b8 108\n.b8 95\n.b8 107\n.b8 101\n.b8 114\n.b8 110\n.b8 101\n.b8 108\n.b8 0\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","matmul_kernel.json":"{\"hash\": \"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 1, \"num_ctas\": 1, \"num_stages\": 1, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.8.0\", \"tensordesc_meta\": [], \"shared\": 1024, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"matmul_kernel\"}","matmul_kernel.sass":"Function:matmul_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 46\n /*0010*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 311\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*0020*/ S2R R7, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 94\n /*0030*/ UMOV UR4, 0x400 ;\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0040*/ ULDC.64 UR6, c[0x0][0x208] ;\n\t//## File \".nv_debug_ptx_txt\", line 94\n /*0050*/ S2UR UR5, SR_CgaCtaId ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*0060*/ VIADD R0, R9, 0xf ;\n\t//## File \".nv_debug_ptx_txt\", line 57\n /*0070*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 94\n /*0080*/ ULEA UR4, UR5, UR4, 0x18 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \".nv_debug_ptx_txt\", line 59\n /*0090*/ LEA.HI R3, R3, R0, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*00a0*/ IABS R10, R7 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \".nv_debug_ptx_txt\", line 60\n /*00b0*/ SHF.R.S32.HI R4, RZ, 0x4, R3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*00c0*/ IABS R5, R4.reuse ;\n /*00d0*/ IABS R12, R4 ;\n /*00e0*/ I2F.RP R0, R5 ;\n /*00f0*/ MUFU.RCP R0, R0 ;\n /*0100*/ VIADD R2, R0, 0xffffffe ;\n /*0110*/ IMAD.MOV R0, RZ, RZ, -R12 ;\n /*0120*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0130*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0140*/ IMAD.MOV R6, RZ, RZ, -R3 ;\n /*0150*/ IMAD R11, R6, R5, RZ ;\n /*0160*/ IMAD.MOV.U32 R6, RZ, RZ, R10 ;\n /*0170*/ IMAD.HI.U32 R3, R3, R11, R2 ;\n /*0180*/ IMAD.HI.U32 R3, R3, R6, RZ ;\n /*0190*/ IMAD R0, R3, R0, R6 ;\n /*01a0*/ ISETP.GT.U32.AND P1, PT, R5, R0, PT ;\n /*01b0*/ @!P1 IMAD.IADD R0, R0, 0x1, -R5 ;\n /*01c0*/ @!P1 VIADD R3, R3, 0x1 ;\n /*01d0*/ ISETP.NE.AND P1, PT, R4, RZ, PT ;\n /*01e0*/ ISETP.GE.U32.AND P0, PT, R0, R5, PT ;\n /*01f0*/ LOP3.LUT R0, R7, R4, RZ, 0x3c, !PT ;\n /*0200*/ ISETP.GE.AND P2, PT, R0, RZ, PT ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 46\n /*0210*/ VIADD R0, R8, 0xf ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0220*/ @P0 VIADD R3, R3, 0x1 ;\n /*0230*/ IMAD.MOV.U32 R12, RZ, RZ, R3 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*0240*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0250*/ @!P2 IMAD.MOV R12, RZ, RZ, -R12 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 50\n /*0260*/ LEA.HI R3, R3, R0, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0270*/ @!P1 LOP3.LUT R12, RZ, R4, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0280*/ LEA.HI.SX32 R3, R3, -R12, 0x1c ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0290*/ IMAD.MOV R5, RZ, RZ, -R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*02a0*/ VIMNMX R15, R3, 0x1, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*02b0*/ IMAD R4, R4, R5, R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*02c0*/ IABS R11, R15.reuse ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*02d0*/ IABS R14, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*02e0*/ I2F.RP R0, R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*02f0*/ LOP3.LUT R4, R4, R15, RZ, 0x3c, !PT ;\n /*0300*/ ISETP.GE.AND P2, PT, R4, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0310*/ MUFU.RCP R0, R0 ;\n /*0320*/ VIADD R2, R0, 0xffffffe ;\n /*0330*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0340*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0350*/ IMAD.MOV R10, RZ, RZ, -R3 ;\n /*0360*/ IMAD R5, R10, R11, RZ ;\n /*0370*/ IABS R10, R15 ;\n /*0380*/ IMAD.HI.U32 R3, R3, R5, R2 ;\n /*0390*/ IMAD.MOV R13, RZ, RZ, -R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*03a0*/ IMAD.MOV.U32 R10, RZ, RZ, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*03b0*/ IMAD.HI.U32 R0, R3, R6, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*03c0*/ IMAD.HI.U32 R5, R3, R10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*03d0*/ IMAD R2, R0, R13.reuse, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*03e0*/ IMAD R3, R5, R13, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*03f0*/ S2R R0, SR_TID.X ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 90\n /*0400*/ LDC R10, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0410*/ ISETP.GT.U32.AND P1, PT, R11.reuse, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0420*/ ISETP.GT.U32.AND P0, PT, R11, R3, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0430*/ @!P1 IMAD.IADD R2, R2, 0x1, -R11.reuse ;\n /*0440*/ ISETP.GE.AND P1, PT, R7, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0450*/ @!P0 IMAD.IADD R3, R3, 0x1, -R11 ;\n /*0460*/ @!P0 VIADD R5, R5, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0470*/ ISETP.GT.U32.AND P4, PT, R11, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0480*/ ISETP.GE.U32.AND P3, PT, R3, R11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0490*/ ISETP.NE.AND P0, PT, R15, RZ, PT ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 90\n /*04a0*/ VIADD R10, R10, 0xf ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*04b0*/ LOP3.LUT R3, RZ, R15, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 80\n /*04c0*/ SHF.R.U32.HI R13, RZ, 0x1, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*04d0*/ @!P4 IMAD.IADD R2, R2, 0x1, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*04e0*/ @P3 VIADD R5, R5, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 80\n /*04f0*/ SGXT.U32 R13, R13, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0500*/ @!P1 IMAD.MOV R2, RZ, RZ, -R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0510*/ @!P2 IMAD.MOV R5, RZ, RZ, -R5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0520*/ SEL R2, R3.reuse, R2, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0530*/ SEL R3, R3, R5, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 93\n /*0540*/ ISETP.GT.AND P0, PT, R10, 0xf, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 71\n /*0550*/ IMAD.IADD R4, R12, 0x1, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0560*/ LOP3.LUT R2, R0, 0x1, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 84\n /*0570*/ IMAD.SHL.U32 R4, R4, 0x10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*0580*/ IMAD.SHL.U32 R14, R2, 0x8, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 85\n /*0590*/ IMAD R12, R3, 0x10, R14 ;\n\t//## File \".nv_debug_ptx_txt\", line 86\n /*05a0*/ LOP3.LUT R3, R4, R13, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 95\n /*05b0*/ @!P0 BRA `(.L_x_0) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*05c0*/ IABS R17, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 337\n\t//## File \".nv_debug_ptx_txt\", line 107\n /*05d0*/ LDC R39, c[0x0][0x238] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*05e0*/ IABS R16, R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*05f0*/ ULDC UR5, c[0x0][0x234] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0600*/ I2F.RP R11, R17 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0610*/ ISETP.GE.AND P2, PT, R12, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0620*/ ULDC.64 UR8, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0630*/ ISETP.NE.AND P4, PT, R9, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0640*/ ULDC.64 UR10, c[0x0][0x210] ;\n\t//## File \".nv_debug_ptx_txt\", line 119\n /*0650*/ IMAD.SHL.U32 R37, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 121\n /*0660*/ LOP3.LUT R38, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0670*/ I2F.RP R15, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 129\n /*0680*/ LOP3.LUT R40, R0, 0x4, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0690*/ MUFU.RCP R11, R11 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*06a0*/ MUFU.RCP R15, R15 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*06b0*/ VIADD R4, R11, 0xffffffe ;\n /*06c0*/ F2I.FTZ.U32.TRUNC.NTZ R5, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*06d0*/ VIADD R6, R15, 0xffffffe ;\n /*06e0*/ IABS R15, R12 ;\n /*06f0*/ F2I.FTZ.U32.TRUNC.NTZ R7, R6 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0700*/ IMAD.MOV.U32 R4, RZ, RZ, RZ ;\n /*0710*/ IMAD.MOV R18, RZ, RZ, -R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0720*/ IMAD.MOV.U32 R6, RZ, RZ, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0730*/ IMAD R19, R18, R17, RZ ;\n /*0740*/ IABS R18, R3 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0750*/ IMAD.MOV R21, RZ, RZ, -R7 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0760*/ IMAD.HI.U32 R4, R5, R19, R4 ;\n /*0770*/ IMAD.MOV.U32 R5, RZ, RZ, R18 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0780*/ IMAD R11, R21, R16, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0790*/ IMAD.HI.U32 R4, R4, R5, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*07a0*/ IMAD.HI.U32 R6, R7, R11, R6 ;\n /*07b0*/ IMAD.MOV.U32 R7, RZ, RZ, R15 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*07c0*/ IMAD.MOV R4, RZ, RZ, -R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*07d0*/ IMAD.HI.U32 R6, R6, R7, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*07e0*/ IMAD R4, R17, R4, R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*07f0*/ IMAD.MOV R6, RZ, RZ, -R6 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0800*/ ISETP.GT.U32.AND P0, PT, R17, R4, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0810*/ IMAD R5, R16, R6, R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 130\n /*0820*/ SHF.R.S32.HI R6, RZ, 0x2, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 127\n /*0830*/ IMAD.SHL.U32 R7, R0, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0840*/ ISETP.GT.U32.AND P1, PT, R16, R5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 130\n /*0850*/ SGXT R6, R6, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0860*/ @!P0 IMAD.IADD R4, R4, 0x1, -R17 ;\n /*0870*/ ISETP.GT.U32.AND P3, PT, R17, R4, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0880*/ @!P1 IMAD.IADD R5, R5, 0x1, -R16 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0890*/ ISETP.GE.AND P1, PT, R3, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*08a0*/ ISETP.GT.U32.AND P0, PT, R16, R5, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*08b0*/ @!P3 IMAD.IADD R4, R4, 0x1, -R17 ;\n /*08c0*/ ISETP.NE.AND P3, PT, R8, RZ, PT ;\n /*08d0*/ IMAD.MOV.U32 R11, RZ, RZ, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*08e0*/ @!P0 IMAD.IADD R5, R5, 0x1, -R16 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*08f0*/ @!P1 IMAD.MOV R11, RZ, RZ, -R11 ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0900*/ IMAD.MOV.U32 R4, RZ, RZ, R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 102\n /*0910*/ SHF.R.S32.HI R5, RZ, 0x1f, R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0920*/ @!P3 LOP3.LUT R11, RZ, R8, RZ, 0x33, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*0930*/ @!P2 IMAD.MOV R4, RZ, RZ, -R4 ;\n /*0940*/ @!P4 LOP3.LUT R4, RZ, R9, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 131\n /*0950*/ LOP3.LUT R9, R6, 0x90, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0960*/ IMAD R6, R11, UR5, R14 ;\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*0970*/ SHF.R.S32.HI R11, RZ, 0x3, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0980*/ IMAD R4, R13, R39, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 133\n /*0990*/ LOP3.LUT R8, R9, 0x10, R0, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 337\n\t//## File \".nv_debug_ptx_txt\", line 107\n /*09a0*/ IMAD.SHL.U32 R39, R39, 0x10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*09b0*/ SGXT R11, R11, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 39\n /*09c0*/ ULDC UR5, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*09d0*/ LEA R36, P1, R6, UR10, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*09e0*/ LEA R24, P0, R4, UR8, 0x1 ;\n\t//## File \".nv_debug_ptx_txt\", line 104\n /*09f0*/ LEA.HI R15, R5, R10, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 123\n /*0a00*/ LOP3.LUT R16, R11, 0x90, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 146\n /*0a10*/ CS2R R10, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 134\n /*0a20*/ LOP3.LUT R35, R8, 0x160, R7, 0xf8, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 146\n /*0a30*/ CS2R R8, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0a40*/ LEA.HI.X.SX32 R41, R6, UR11, 0x1, P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 146\n /*0a50*/ CS2R R6, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0a60*/ LEA.HI.X.SX32 R25, R4, UR9, 0x1, P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 146\n /*0a70*/ CS2R R4, SRZ ;\n\t//## File \".nv_debug_ptx_txt\", line 105\n /*0a80*/ SHF.R.S32.HI R15, RZ, 0x4, R15 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0a90*/ LOP3.LUT R34, R16, 0x170, R37, 0x78, !PT ;\n.L_x_1:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 149\n /*0aa0*/ ISETP.GE.AND P0, PT, R14, UR5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 161\n /*0ab0*/ CS2R R16, SRZ ;\n /*0ac0*/ CS2R R18, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0ad0*/ ISETP.GE.AND P1, PT, R13, UR5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 182\n /*0ae0*/ CS2R R20, SRZ ;\n /*0af0*/ CS2R R22, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 161\n /*0b00*/ @!P0 IMAD.MOV.U32 R32, RZ, RZ, R36 ;\n /*0b10*/ @!P0 IMAD.MOV.U32 R33, RZ, RZ, R41 ;\n /*0b20*/ @!P0 LDG.E.128 R16, desc[UR6][R32.64] ;\n\t//## File \".nv_debug_ptx_txt\", line 164\n /*0b30*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n /*0b40*/ IMAD.MOV.U32 R32, RZ, RZ, R24 ;\n /*0b50*/ IMAD.MOV.U32 R33, RZ, RZ, R25 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 182\n /*0b60*/ @!P1 LDG.E.128 R20, desc[UR6][R32.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 210\n /*0b70*/ VIADD R15, R15, 0xffffffff ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 207\n /*0b80*/ IADD3 R36, P1, R36, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 211\n /*0b90*/ UIADD3 UR5, UR5, -0x10, URZ ;\n\t//## File \".nv_debug_ptx_txt\", line 212\n /*0ba0*/ ISETP.NE.U32.AND P0, PT, R15, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 207\n /*0bb0*/ IMAD.X R41, RZ, RZ, R41, P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 329\n\t//## File \".nv_debug_ptx_txt\", line 167\n /*0bc0*/ STS.128 [R34+UR4], R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 187\n /*0bd0*/ STS.128 [R34+UR4+0x200], R20 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 329\n\t//## File \".nv_debug_ptx_txt\", line 191\n /*0be0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 192\n /*0bf0*/ LDSM.16.M88.4 R24, [R35+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 194\n /*0c00*/ LDSM.16.MT88.4 R28, [R35+UR4+0x200] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 198\n /*0c10*/ HMMA.16816.F32 R4, R24.reuse, R28, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 203\n /*0c20*/ HMMA.16816.F32 R8, R24, R30, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 209\n /*0c30*/ IMAD.WIDE R24, R39, 0x2, R32 ;\n /*0c40*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 213\n /*0c50*/ @P0 BRA `(.L_x_1) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 338\n\t//## File \".nv_debug_ptx_txt\", line 216\n /*0c60*/ F2FP.F16.F32.PACK_AB R11, R11, R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 217\n /*0c70*/ F2FP.F16.F32.PACK_AB R10, R9, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 218\n /*0c80*/ F2FP.F16.F32.PACK_AB R9, R7, R6 ;\n\t//## File \".nv_debug_ptx_txt\", line 219\n /*0c90*/ F2FP.F16.F32.PACK_AB R8, R5, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 220\n /*0ca0*/ BRA `(.L_x_2) ;\n.L_x_0:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 223\n /*0cb0*/ IMAD.SHL.U32 R37, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 224\n /*0cc0*/ LOP3.LUT R40, R0.reuse, 0x4, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 229\n /*0cd0*/ CS2R R8, SRZ ;\n\t//## File \".nv_debug_ptx_txt\", line 225\n /*0ce0*/ LOP3.LUT R38, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 229\n /*0cf0*/ CS2R R10, SRZ ;\n.L_x_2:\n\t//## File \".nv_debug_ptx_txt\", line 243\n /*0d00*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 245\n /*0d10*/ ISETP.NE.U32.AND P0, PT, R40, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 254\n /*0d20*/ IMAD.SHL.U32 R5, R0, 0x8, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 257\n /*0d30*/ ISETP.NE.U32.AND P1, PT, R38, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 247\n /*0d40*/ IMAD.SHL.U32 R0, R0, 0x4, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 246\n /*0d50*/ SEL R4, RZ, 0x140, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 239\n /*0d60*/ ULDC.64 UR10, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 255\n /*0d70*/ LOP3.LUT R5, R5, 0xb0, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 232\n /*0d80*/ ULDC UR5, c[0x0][0x23c] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 249\n /*0d90*/ LOP3.LUT R37, R4, 0xb0, R37, 0xf8, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0da0*/ ULDC.64 UR8, c[0x0][0x220] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 239\n /*0db0*/ ISETP.GE.AND P0, PT, R12, UR11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 259\n /*0dc0*/ IMAD R13, R2, 0x40, R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 250\n /*0dd0*/ LOP3.LUT R37, R37, 0x40, R0, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 232\n /*0de0*/ IMAD R0, R3.reuse, UR5, RZ ;\n /*0df0*/ ISETP.LT.AND P0, PT, R3, UR10, !P0 ;\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0e00*/ IMAD.SHL.U32 R5, R12, 0x2, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 258\n /*0e10*/ SEL R4, RZ, 0x140, !P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 234\n /*0e20*/ IMAD.SHL.U32 R2, R0, 0x2, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 251\n /*0e30*/ VIADD R37, R37, UR4 ;\n\t//## File \".nv_debug_ptx_txt\", line 260\n /*0e40*/ LOP3.LUT R4, R13, R4, RZ, 0x3c, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0e50*/ IMAD.HI R7, R12, 0x2, RZ ;\n /*0e60*/ IADD3 R2, P1, P2, R5, UR8, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 252\n /*0e70*/ STSM.16.M88.4 [R37], R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 234\n /*0e80*/ IMAD.HI R0, R0, 0x2, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0e90*/ IADD3.X R3, R7, UR9, R0, P1, P2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 253\n /*0ea0*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 264\n /*0eb0*/ @!P0 EXIT ;\n\t//## File \".nv_debug_ptx_txt\", line 262\n /*0ec0*/ LDS.128 R4, [R4+UR4] ;\n\t//## File \".nv_debug_ptx_txt\", line 264\n /*0ed0*/ STG.E.128 desc[UR6][R2.64], R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 268\n /*0ee0*/ EXIT ;\n.L_x_3:\n /*0ef0*/ BRA `(.L_x_3);\n /*0f00*/ NOP;\n /*0f10*/ NOP;\n /*0f20*/ NOP;\n /*0f30*/ NOP;\n /*0f40*/ NOP;\n /*0f50*/ NOP;\n /*0f60*/ NOP;\n /*0f70*/ NOP;\n.L_x_4:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":266,"end_line":344,"code":" @triton.autotune(\n configs=[\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 16,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 32,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n ],\n key=[\"M\", \"N\", \"K\"],\n )\n @triton.jit\n def matmul_kernel(\n a,\n b,\n c,\n M,\n N,\n K,\n stride_am,\n stride_ak,\n stride_bk,\n stride_bn,\n stride_cm,\n stride_cn,\n BLOCK_SIZE_M: tl.constexpr,\n BLOCK_SIZE_N: tl.constexpr,\n BLOCK_SIZE_K: tl.constexpr,\n GROUP_SIZE_M: tl.constexpr,\n ACTIVATION: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)\n num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)\n num_pid_in_group = GROUP_SIZE_M * num_pid_n\n group_id = pid // num_pid_in_group\n first_pid_m = group_id * GROUP_SIZE_M\n group_size = min(num_pid_m - first_pid_m, GROUP_SIZE_M)\n pid_m = first_pid_m + (pid % group_size)\n pid_n = (pid % num_pid_in_group) // group_size\n\n offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M\n offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N\n offs_k = tl.arange(0, BLOCK_SIZE_K)\n a_ptrs = a + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)\n b_ptrs = b + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)\n\n accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)\n for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):\n a_block = tl.load(\n a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0\n )\n b_block = tl.load(\n b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0\n )\n accumulator += tl.dot(a_block, b_block)\n a_ptrs += BLOCK_SIZE_K * stride_ak\n b_ptrs += BLOCK_SIZE_K * stride_bk\n c_block = accumulator.to(tl.float16)\n\n offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)\n offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)\n c_ptrs = c + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]\n c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)\n tl.store(c_ptrs, c_block, mask=c_mask)\n"}}} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":272,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":283,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":283,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":283,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":283,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":283,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":295,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":308,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":308,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":308,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"launch","pid":4104992,"name":"matmul_kernel","function":94159363288336,"stream":0,"grid":[1],"compilation_metadata":{"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2b5b149e9d5ae5161f98984882ad1110dbfd0ae3001b0f103baae7ec6427464c","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1024,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32},"extracted_args":{"a":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00000"},"b":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00200"},"c":{"type":"tensor","shape":[16,16],"dtype":"torch.float16","device":"cuda:0","stride":[16,1],"numel":256,"is_contiguous":true,"element_size":2,"storage_offset":0,"memory_usage":512,"data_ptr":"0x7fc747e00400"},"M":{"type":"int","value":16},"N":{"type":"int","value":16},"K":{"type":"int","value":16},"stride_am":{"type":"int","value":16},"stride_ak":{"type":"int","value":1},"stride_bk":{"type":"int","value":16},"stride_bn":{"type":"int","value":1},"stride_cm":{"type":"int","value":16},"stride_cn":{"type":"int","value":1},"BLOCK_SIZE_M":{"type":"int","value":16},"BLOCK_SIZE_N":{"type":"int","value":16},"BLOCK_SIZE_K":{"type":"int","value":16},"GROUP_SIZE_M":{"type":"int","value":1},"ACTIVATION":{"type":"NoneType","repr":"None"}},"extracted_inductor_args":{},"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":308,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":779,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel.run(grid_0, grid_1, grid_2, stream, kernel.function, kernel.packed_metadata, launch_metadata,"},{"line":331,"name":"__call__","filename":"/python/lib/python3.14/site-packages/triton/backends/nvidia/driver.py","loc":"self.launch(gridX, gridY, gridZ, stream, function, self.launch_cooperative_grid, self.launch_pdl,"},{"line":2027,"name":"__call__","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"launch\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"compilation","pid":4104992,"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":251,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"benchmark()"},{"line":240,"name":"benchmark","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"timings = {config: self._bench(*args, config=config, **kwargs) for config in pruned_configs}"},{"line":169,"name":"_bench","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"return self.do_bench(kernel_call, quantiles=(0.5, 0.2, 0.8))"},{"line":272,"name":"do_bench","filename":"/python/lib/python3.14/site-packages/triton/testing.py","loc":"fn()"},{"line":155,"name":"kernel_call","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"self.fn.run("},{"line":757,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":902,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"2645961e93d089721e51c9760195f3a7cabfac6742babb6be3a4dbc4331529eb","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"matmul_kernel","num_ctas":1,"num_stages":1,"num_warps":1,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":1536,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[["tt.divisibility",16]],"(5,)":[["tt.divisibility",16]],"(6,)":[["tt.divisibility",16]],"(7,)":[],"(8,)":[["tt.divisibility",16]],"(9,)":[],"(10,)":[["tt.divisibility",16]],"(11,)":[],"(12,)":[],"(13,)":[],"(14,)":[],"(15,)":[],"(16,)":[]},"src_cache_key":"528f8911847b1f39f462157322ec9a674a5565b9b74d621d5ff09625eb710f38","src_constants":{"(7,)":1,"(9,)":1,"(11,)":1,"(12,)":32,"(13,)":16,"(14,)":16,"(15,)":1,"(16,)":null},"times":{"ir_initialization":1079,"lowering_stages":[],"store_results":0}},"file_path":{"matmul_kernel.source":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.source","matmul_kernel.ttir":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.ttir","matmul_kernel.ttgir":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.ttgir","matmul_kernel.llir":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.llir","matmul_kernel.ptx":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.ptx","matmul_kernel.cubin":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.cubin","matmul_kernel.json":"/home/user/.triton/cache/EZCZMHUT2CEXEHSRZF3ADFPTU7FL7LDHIK5LW27DUTN4IMYVFHVQ/matmul_kernel.json"},"file_content":{"matmul_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:5)\n#loc47 = loc(\"a\"(#loc))\n#loc48 = loc(\"b\"(#loc))\n#loc49 = loc(\"c\"(#loc))\n#loc50 = loc(\"M\"(#loc))\n#loc51 = loc(\"N\"(#loc))\n#loc52 = loc(\"K\"(#loc))\n#loc53 = loc(\"stride_am\"(#loc))\n#loc54 = loc(\"stride_bk\"(#loc))\n#loc55 = loc(\"stride_cm\"(#loc))\nmodule {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %num_pid_m = arith.constant 31 : i32 loc(#loc99)\n %cst = arith.constant dense<0.000000e+00> : tensor<16x16xf16> loc(#loc1)\n %cst_0 = arith.constant dense<0.000000e+00> : tensor<32x16xf16> loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc3)\n %cst_1 = arith.constant dense<16> : tensor<32x16xi32> loc(#loc1)\n %cst_2 = arith.constant dense<0.000000e+00> : tensor<32x16xf32> loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c32_i32 = arith.constant 32 : i32 loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc57)\n %num_pid_m_3 = arith.addi %M, %num_pid_m : i32 loc(#loc100)\n %num_pid_m_4 = arith.divsi %num_pid_m_3, %c32_i32 : i32 loc(#loc101)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc102)\n %num_pid_n_5 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc103)\n %group_id = arith.divsi %pid, %num_pid_n_5 : i32 loc(#loc59)\n %group_size = arith.subi %num_pid_m_4, %group_id : i32 loc(#loc60)\n %group_size_6 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc61)\n %pid_m = arith.remsi %pid, %group_size_6 : i32 loc(#loc62)\n %pid_m_7 = arith.addi %group_id, %pid_m : i32 loc(#loc63)\n %pid_n = arith.remsi %pid, %num_pid_n_5 : i32 loc(#loc64)\n %pid_n_8 = arith.divsi %pid_n, %group_size_6 : i32 loc(#loc65)\n %offs_am = arith.muli %pid_m_7, %c32_i32 : i32 loc(#loc66)\n %offs_am_9 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32> loc(#loc67)\n %offs_am_10 = tt.splat %offs_am : i32 -> tensor<32xi32> loc(#loc66)\n %offs_am_11 = arith.addi %offs_am_10, %offs_am_9 : tensor<32xi32> loc(#loc66)\n %offs_am_12 = tt.splat %M : i32 -> tensor<32xi32> loc(#loc68)\n %offs_am_13 = arith.remsi %offs_am_11, %offs_am_12 : tensor<32xi32> loc(#loc68)\n %offs_bn = arith.muli %pid_n_8, %c16_i32 : i32 loc(#loc69)\n %offs_bn_14 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32> loc(#loc70)\n %offs_bn_15 = tt.splat %offs_bn : i32 -> tensor<16xi32> loc(#loc69)\n %offs_bn_16 = arith.addi %offs_bn_15, %offs_bn_14 : tensor<16xi32> loc(#loc69)\n %offs_bn_17 = tt.splat %N : i32 -> tensor<16xi32> loc(#loc71)\n %offs_bn_18 = arith.remsi %offs_bn_16, %offs_bn_17 : tensor<16xi32> loc(#loc71)\n %a_ptrs = tt.expand_dims %offs_am_13 {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc72)\n %a_ptrs_19 = tt.splat %stride_am : i32 -> tensor<32x1xi32> loc(#loc72)\n %a_ptrs_20 = arith.muli %a_ptrs, %a_ptrs_19 : tensor<32x1xi32> loc(#loc72)\n %a_ptrs_21 = tt.expand_dims %offs_bn_14 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc73)\n %a_ptrs_22 = tt.broadcast %a_ptrs_20 : tensor<32x1xi32> -> tensor<32x16xi32> loc(#loc72)\n %a_ptrs_23 = tt.broadcast %a_ptrs_21 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc72)\n %a_ptrs_24 = arith.addi %a_ptrs_22, %a_ptrs_23 : tensor<32x16xi32> loc(#loc72)\n %a_ptrs_25 = tt.splat %a : !tt.ptr -> tensor<32x16x!tt.ptr> loc(#loc74)\n %a_ptrs_26 = tt.addptr %a_ptrs_25, %a_ptrs_24 : tensor<32x16x!tt.ptr>, tensor<32x16xi32> loc(#loc74)\n %b_ptrs = tt.expand_dims %offs_bn_14 {axis = 1 : i32} : tensor<16xi32> -> tensor<16x1xi32> loc(#loc75)\n %b_ptrs_27 = tt.splat %stride_bk : i32 -> tensor<16x1xi32> loc(#loc75)\n %b_ptrs_28 = arith.muli %b_ptrs, %b_ptrs_27 : tensor<16x1xi32> loc(#loc75)\n %b_ptrs_29 = tt.expand_dims %offs_bn_18 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc76)\n %b_ptrs_30 = tt.broadcast %b_ptrs_28 : tensor<16x1xi32> -> tensor<16x16xi32> loc(#loc75)\n %b_ptrs_31 = tt.broadcast %b_ptrs_29 : tensor<1x16xi32> -> tensor<16x16xi32> loc(#loc75)\n %b_ptrs_32 = arith.addi %b_ptrs_30, %b_ptrs_31 : tensor<16x16xi32> loc(#loc75)\n %b_ptrs_33 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr> loc(#loc77)\n %b_ptrs_34 = tt.addptr %b_ptrs_33, %b_ptrs_32 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc77)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc78)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc79)\n %accumulator:3 = scf.for %k = %c0_i32 to %1 step %c1_i32 iter_args(%a_ptrs_49 = %a_ptrs_26, %b_ptrs_50 = %b_ptrs_34, %accumulator_51 = %cst_2) -> (tensor<32x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<32x16xf32>) : i32 {\n %a_block = arith.muli %k, %c16_i32 : i32 loc(#loc81)\n %a_block_52 = arith.subi %K, %a_block : i32 loc(#loc82)\n %a_block_53 = tt.splat %a_block_52 : i32 -> tensor<1x16xi32> loc(#loc83)\n %a_block_54 = arith.cmpi slt, %a_ptrs_21, %a_block_53 : tensor<1x16xi32> loc(#loc83)\n %a_block_55 = tt.broadcast %a_block_54 : tensor<1x16xi1> -> tensor<32x16xi1> loc(#loc84)\n %a_block_56 = tt.load %a_ptrs_49, %a_block_55, %cst_0 : tensor<32x16x!tt.ptr> loc(#loc84)\n %b_block = tt.splat %a_block_52 : i32 -> tensor<16x1xi32> loc(#loc85)\n %b_block_57 = arith.cmpi slt, %b_ptrs, %b_block : tensor<16x1xi32> loc(#loc85)\n %b_block_58 = tt.broadcast %b_block_57 : tensor<16x1xi1> -> tensor<16x16xi1> loc(#loc86)\n %b_block_59 = tt.load %b_ptrs_50, %b_block_58, %cst : tensor<16x16x!tt.ptr> loc(#loc86)\n %accumulator_60 = tt.dot %a_block_56, %b_block_59, %accumulator_51, inputPrecision = tf32 : tensor<32x16xf16> * tensor<16x16xf16> -> tensor<32x16xf32> loc(#loc87)\n %a_ptrs_61 = tt.addptr %a_ptrs_49, %cst_1 : tensor<32x16x!tt.ptr>, tensor<32x16xi32> loc(#loc88)\n %b_ptrs_62 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc89)\n %b_ptrs_63 = tt.splat %b_ptrs_62 : i32 -> tensor<16x16xi32> loc(#loc90)\n %b_ptrs_64 = tt.addptr %b_ptrs_50, %b_ptrs_63 : tensor<16x16x!tt.ptr>, tensor<16x16xi32> loc(#loc90)\n scf.yield %a_ptrs_61, %b_ptrs_64, %accumulator_60 : tensor<32x16x!tt.ptr>, tensor<16x16x!tt.ptr>, tensor<32x16xf32> loc(#loc3)\n } loc(#loc105)\n %c_block = arith.truncf %accumulator#2 : tensor<32x16xf32> to tensor<32x16xf16> loc(#loc91)\n %c_ptrs = tt.expand_dims %offs_am_11 {axis = 1 : i32} : tensor<32xi32> -> tensor<32x1xi32> loc(#loc92)\n %c_ptrs_35 = tt.splat %stride_cm : i32 -> tensor<32x1xi32> loc(#loc93)\n %c_ptrs_36 = arith.muli %c_ptrs_35, %c_ptrs : tensor<32x1xi32> loc(#loc93)\n %c_ptrs_37 = tt.splat %c : !tt.ptr -> tensor<32x1x!tt.ptr> loc(#loc94)\n %c_ptrs_38 = tt.addptr %c_ptrs_37, %c_ptrs_36 : tensor<32x1x!tt.ptr>, tensor<32x1xi32> loc(#loc94)\n %c_ptrs_39 = tt.expand_dims %offs_bn_16 {axis = 0 : i32} : tensor<16xi32> -> tensor<1x16xi32> loc(#loc95)\n %c_ptrs_40 = tt.broadcast %c_ptrs_38 : tensor<32x1x!tt.ptr> -> tensor<32x16x!tt.ptr> loc(#loc94)\n %c_ptrs_41 = tt.broadcast %c_ptrs_39 : tensor<1x16xi32> -> tensor<32x16xi32> loc(#loc94)\n %c_ptrs_42 = tt.addptr %c_ptrs_40, %c_ptrs_41 : tensor<32x16x!tt.ptr>, tensor<32x16xi32> loc(#loc94)\n %c_mask = tt.splat %M : i32 -> tensor<32x1xi32> loc(#loc96)\n %c_mask_43 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<32x1xi32> loc(#loc96)\n %c_mask_44 = tt.splat %N : i32 -> tensor<1x16xi32> loc(#loc97)\n %c_mask_45 = arith.cmpi slt, %c_ptrs_39, %c_mask_44 : tensor<1x16xi32> loc(#loc97)\n %c_mask_46 = tt.broadcast %c_mask_43 : tensor<32x1xi1> -> tensor<32x16xi1> loc(#loc98)\n %c_mask_47 = tt.broadcast %c_mask_45 : tensor<1x16xi1> -> tensor<32x16xi1> loc(#loc98)\n %c_mask_48 = arith.andi %c_mask_46, %c_mask_47 : tensor<32x16xi1> loc(#loc98)\n tt.store %c_ptrs_42, %c_block, %c_mask_48 : tensor<32x16x!tt.ptr> loc(#loc46)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:21)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:9)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:15)\n#loc5 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:13)\n#loc6 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:12)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:21)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:20)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:26)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:22)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:32)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:17)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:18)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:17)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:20)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:43)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:19)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:20)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:43)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:19)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:23)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:54)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:23)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:53)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:27)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:52)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:48)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:30)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":329:23)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:30)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":332:23)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:28)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:13)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:23)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:13)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:19)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:34)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:22)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:18)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:65)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:19)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:44)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:18)\n#loc46 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:9)\n#loc56 = loc(\"num_pid_m\"(#loc2))\n#loc57 = loc(\"pid\"(#loc4))\n#loc58 = loc(\"num_pid_n\"(#loc7))\n#loc59 = loc(\"group_id\"(#loc8))\n#loc60 = loc(\"group_size\"(#loc9))\n#loc61 = loc(\"group_size\"(#loc10))\n#loc62 = loc(\"pid_m\"(#loc11))\n#loc63 = loc(\"pid_m\"(#loc12))\n#loc64 = loc(\"pid_n\"(#loc13))\n#loc65 = loc(\"pid_n\"(#loc14))\n#loc66 = loc(\"offs_am\"(#loc15))\n#loc67 = loc(\"offs_am\"(#loc16))\n#loc68 = loc(\"offs_am\"(#loc17))\n#loc69 = loc(\"offs_bn\"(#loc18))\n#loc70 = loc(\"offs_bn\"(#loc19))\n#loc71 = loc(\"offs_bn\"(#loc20))\n#loc72 = loc(\"a_ptrs\"(#loc21))\n#loc73 = loc(\"a_ptrs\"(#loc22))\n#loc74 = loc(\"a_ptrs\"(#loc23))\n#loc75 = loc(\"b_ptrs\"(#loc24))\n#loc76 = loc(\"b_ptrs\"(#loc25))\n#loc77 = loc(\"b_ptrs\"(#loc26))\n#loc78 = loc(callsite(#loc5 at #loc27))\n#loc79 = loc(callsite(#loc6 at #loc27))\n#loc80 = loc(\"a_ptrs\"(#loc3))\n#loc81 = loc(\"a_block\"(#loc28))\n#loc82 = loc(\"a_block\"(#loc29))\n#loc83 = loc(\"a_block\"(#loc30))\n#loc84 = loc(\"a_block\"(#loc31))\n#loc85 = loc(\"b_block\"(#loc32))\n#loc86 = loc(\"b_block\"(#loc33))\n#loc87 = loc(\"accumulator\"(#loc34))\n#loc88 = loc(\"a_ptrs\"(#loc35))\n#loc89 = loc(\"b_ptrs\"(#loc36))\n#loc90 = loc(\"b_ptrs\"(#loc37))\n#loc91 = loc(\"c_block\"(#loc38))\n#loc92 = loc(\"c_ptrs\"(#loc39))\n#loc93 = loc(\"c_ptrs\"(#loc40))\n#loc94 = loc(\"c_ptrs\"(#loc41))\n#loc95 = loc(\"c_ptrs\"(#loc42))\n#loc96 = loc(\"c_mask\"(#loc43))\n#loc97 = loc(\"c_mask\"(#loc44))\n#loc98 = loc(\"c_mask\"(#loc45))\n#loc99 = loc(callsite(#loc1 at #loc56))\n#loc100 = loc(callsite(#loc5 at #loc56))\n#loc101 = loc(callsite(#loc6 at #loc56))\n#loc102 = loc(callsite(#loc5 at #loc58))\n#loc103 = loc(callsite(#loc6 at #loc58))\n#loc104 = loc(\"b_ptrs\"(#loc80))\n#loc105 = loc(\"accumulator\"(#loc104))\n","matmul_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1, 8], threadsPerWarp = [16, 2], warpsPerCTA = [1, 1], order = [1, 0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":292:5)\n#mma = #ttg.nvidia_mma<{versionMajor = 2, versionMinor = 0, warpsPerCTA = [1, 1], instrShape = [16, 8]}>\n#shared = #ttg.swizzled_shared<{vec = 8, perPhase = 4, maxPhase = 2, order = [1, 0]}>\n#smem = #ttg.shared_memory\n#loc47 = loc(\"a\"(#loc))\n#loc48 = loc(\"b\"(#loc))\n#loc49 = loc(\"c\"(#loc))\n#loc50 = loc(\"M\"(#loc))\n#loc51 = loc(\"N\"(#loc))\n#loc52 = loc(\"K\"(#loc))\n#loc53 = loc(\"stride_am\"(#loc))\n#loc54 = loc(\"stride_bk\"(#loc))\n#loc55 = loc(\"stride_cm\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 1 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @matmul_kernel(%a: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a\"(#loc)), %b: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b\"(#loc)), %c: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c\"(#loc)), %M: i32 {tt.divisibility = 16 : i32} loc(\"M\"(#loc)), %N: i32 {tt.divisibility = 16 : i32} loc(\"N\"(#loc)), %K: i32 {tt.divisibility = 16 : i32} loc(\"K\"(#loc)), %stride_am: i32 {tt.divisibility = 16 : i32} loc(\"stride_am\"(#loc)), %stride_bk: i32 {tt.divisibility = 16 : i32} loc(\"stride_bk\"(#loc)), %stride_cm: i32 {tt.divisibility = 16 : i32} loc(\"stride_cm\"(#loc))) attributes {noinline = false} {\n %cst = arith.constant dense<0.000000e+00> : tensor<32x16xf32, #mma> loc(#loc1)\n %cst_0 = arith.constant dense<0.000000e+00> : tensor<16x16xf16, #blocked> loc(#loc1)\n %cst_1 = arith.constant dense<0.000000e+00> : tensor<32x16xf16, #blocked> loc(#loc1)\n %c1_i32 = arith.constant 1 : i32 loc(#loc1)\n %c32_i32 = arith.constant 32 : i32 loc(#loc1)\n %c16_i32 = arith.constant 16 : i32 loc(#loc1)\n %c15_i32 = arith.constant 15 : i32 loc(#loc1)\n %c31_i32 = arith.constant 31 : i32 loc(#loc1)\n %c0_i32 = arith.constant 0 : i32 loc(#loc1)\n %cst_2 = arith.constant dense<16> : tensor<32x16xi32, #blocked> loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc56)\n %num_pid_m = arith.addi %M, %c31_i32 : i32 loc(#loc99)\n %num_pid_m_3 = arith.divsi %num_pid_m, %c32_i32 : i32 loc(#loc100)\n %num_pid_n = arith.addi %N, %c15_i32 : i32 loc(#loc101)\n %num_pid_n_4 = arith.divsi %num_pid_n, %c16_i32 : i32 loc(#loc102)\n %group_id = arith.divsi %pid, %num_pid_n_4 : i32 loc(#loc59)\n %group_size = arith.subi %num_pid_m_3, %group_id : i32 loc(#loc60)\n %group_size_5 = arith.minsi %group_size, %c1_i32 : i32 loc(#loc61)\n %pid_m = arith.remsi %pid, %group_size_5 : i32 loc(#loc62)\n %pid_m_6 = arith.addi %group_id, %pid_m : i32 loc(#loc63)\n %pid_n = arith.remsi %pid, %num_pid_n_4 : i32 loc(#loc64)\n %pid_n_7 = arith.divsi %pid_n, %group_size_5 : i32 loc(#loc65)\n %offs_am = arith.muli %pid_m_6, %c32_i32 : i32 loc(#loc66)\n %offs_am_8 = tt.make_range {end = 32 : i32, start = 0 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc67)\n %offs_am_9 = tt.splat %offs_am : i32 -> tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc66)\n %offs_am_10 = arith.addi %offs_am_9, %offs_am_8 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc66)\n %offs_am_11 = tt.splat %M : i32 -> tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc68)\n %offs_am_12 = arith.remsi %offs_am_10, %offs_am_11 : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc68)\n %offs_bn = arith.muli %pid_n_7, %c16_i32 : i32 loc(#loc69)\n %offs_bn_13 = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc70)\n %offs_bn_14 = tt.splat %offs_bn : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc69)\n %offs_bn_15 = arith.addi %offs_bn_14, %offs_bn_13 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc69)\n %offs_bn_16 = tt.splat %N : i32 -> tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc71)\n %offs_bn_17 = arith.remsi %offs_bn_15, %offs_bn_16 : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> loc(#loc71)\n %a_ptrs = tt.expand_dims %offs_am_12 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked> loc(#loc72)\n %a_ptrs_18 = tt.splat %stride_am : i32 -> tensor<32x1xi32, #blocked> loc(#loc72)\n %a_ptrs_19 = arith.muli %a_ptrs, %a_ptrs_18 : tensor<32x1xi32, #blocked> loc(#loc72)\n %a_ptrs_20 = tt.expand_dims %offs_bn_13 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc73)\n %a_ptrs_21 = tt.broadcast %a_ptrs_19 : tensor<32x1xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc72)\n %a_ptrs_22 = tt.broadcast %a_ptrs_20 : tensor<1x16xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc72)\n %a_ptrs_23 = arith.addi %a_ptrs_21, %a_ptrs_22 : tensor<32x16xi32, #blocked> loc(#loc72)\n %a_ptrs_24 = tt.splat %a : !tt.ptr -> tensor<32x16x!tt.ptr, #blocked> loc(#loc74)\n %a_ptrs_25 = tt.addptr %a_ptrs_24, %a_ptrs_23 : tensor<32x16x!tt.ptr, #blocked>, tensor<32x16xi32, #blocked> loc(#loc74)\n %b_ptrs = tt.make_range {end = 16 : i32, start = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> loc(#loc75)\n %b_ptrs_26 = tt.expand_dims %b_ptrs {axis = 1 : i32} : tensor<16xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<16x1xi32, #blocked> loc(#loc75)\n %b_ptrs_27 = tt.splat %stride_bk : i32 -> tensor<16x1xi32, #blocked> loc(#loc75)\n %b_ptrs_28 = arith.muli %b_ptrs_26, %b_ptrs_27 : tensor<16x1xi32, #blocked> loc(#loc75)\n %b_ptrs_29 = tt.expand_dims %offs_bn_17 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc76)\n %b_ptrs_30 = tt.broadcast %b_ptrs_28 : tensor<16x1xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc75)\n %b_ptrs_31 = tt.broadcast %b_ptrs_29 : tensor<1x16xi32, #blocked> -> tensor<16x16xi32, #blocked> loc(#loc75)\n %b_ptrs_32 = arith.addi %b_ptrs_30, %b_ptrs_31 : tensor<16x16xi32, #blocked> loc(#loc75)\n %b_ptrs_33 = tt.splat %b : !tt.ptr -> tensor<16x16x!tt.ptr, #blocked> loc(#loc77)\n %b_ptrs_34 = tt.addptr %b_ptrs_33, %b_ptrs_32 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc77)\n %0 = arith.addi %K, %c15_i32 : i32 loc(#loc78)\n %1 = arith.divsi %0, %c16_i32 : i32 loc(#loc79)\n %b_ptrs_35 = arith.muli %stride_bk, %c16_i32 : i32 loc(#loc80)\n %b_ptrs_36 = tt.splat %b_ptrs_35 : i32 -> tensor<16x16xi32, #blocked> loc(#loc81)\n %accumulator:3 = scf.for %k = %c0_i32 to %1 step %c1_i32 iter_args(%accumulator_51 = %cst, %a_ptrs_52 = %a_ptrs_25, %b_ptrs_53 = %b_ptrs_34) -> (tensor<32x16xf32, #mma>, tensor<32x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked>) : i32 {\n %a_block = arith.muli %k, %c16_i32 : i32 loc(#loc83)\n %a_block_54 = arith.subi %K, %a_block : i32 loc(#loc84)\n %a_block_55 = tt.splat %a_block_54 : i32 -> tensor<1x16xi32, #blocked> loc(#loc85)\n %a_block_56 = arith.cmpi slt, %a_ptrs_20, %a_block_55 : tensor<1x16xi32, #blocked> loc(#loc85)\n %a_block_57 = tt.broadcast %a_block_56 : tensor<1x16xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc86)\n %a_block_58 = tt.load %a_ptrs_52, %a_block_57, %cst_1 : tensor<32x16x!tt.ptr, #blocked> loc(#loc86)\n %a_block_59 = ttg.local_alloc %a_block_58 : (tensor<32x16xf16, #blocked>) -> !ttg.memdesc<32x16xf16, #shared, #smem> loc(#loc86)\n %b_block = tt.splat %a_block_54 : i32 -> tensor<16x1xi32, #blocked> loc(#loc87)\n %b_block_60 = arith.cmpi slt, %b_ptrs_26, %b_block : tensor<16x1xi32, #blocked> loc(#loc87)\n %b_block_61 = tt.broadcast %b_block_60 : tensor<16x1xi1, #blocked> -> tensor<16x16xi1, #blocked> loc(#loc88)\n %b_block_62 = tt.load %b_ptrs_53, %b_block_61, %cst_0 : tensor<16x16x!tt.ptr, #blocked> loc(#loc88)\n %b_block_63 = ttg.local_alloc %b_block_62 : (tensor<16x16xf16, #blocked>) -> !ttg.memdesc<16x16xf16, #shared, #smem> loc(#loc88)\n %a_block_64 = ttg.local_load %a_block_59 : !ttg.memdesc<32x16xf16, #shared, #smem> -> tensor<32x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> loc(#loc86)\n %b_block_65 = ttg.local_load %b_block_63 : !ttg.memdesc<16x16xf16, #shared, #smem> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> loc(#loc88)\n %accumulator_66 = tt.dot %a_block_64, %b_block_65, %accumulator_51, inputPrecision = tf32 : tensor<32x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 2}>> * tensor<16x16xf16, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 2}>> -> tensor<32x16xf32, #mma> loc(#loc89)\n %a_ptrs_67 = tt.addptr %a_ptrs_52, %cst_2 : tensor<32x16x!tt.ptr, #blocked>, tensor<32x16xi32, #blocked> loc(#loc90)\n %b_ptrs_68 = tt.addptr %b_ptrs_53, %b_ptrs_36 : tensor<16x16x!tt.ptr, #blocked>, tensor<16x16xi32, #blocked> loc(#loc81)\n scf.yield %accumulator_66, %a_ptrs_67, %b_ptrs_68 : tensor<32x16xf32, #mma>, tensor<32x16x!tt.ptr, #blocked>, tensor<16x16x!tt.ptr, #blocked> loc(#loc29)\n } loc(#loc104)\n %c_block = arith.truncf %accumulator#0 : tensor<32x16xf32, #mma> to tensor<32x16xf16, #mma> loc(#loc91)\n %c_ptrs = tt.expand_dims %offs_am_10 {axis = 1 : i32} : tensor<32xi32, #ttg.slice<{dim = 1, parent = #blocked}>> -> tensor<32x1xi32, #blocked> loc(#loc92)\n %c_ptrs_37 = tt.splat %stride_cm : i32 -> tensor<32x1xi32, #blocked> loc(#loc93)\n %c_ptrs_38 = arith.muli %c_ptrs_37, %c_ptrs : tensor<32x1xi32, #blocked> loc(#loc93)\n %c_ptrs_39 = tt.splat %c : !tt.ptr -> tensor<32x1x!tt.ptr, #blocked> loc(#loc94)\n %c_ptrs_40 = tt.addptr %c_ptrs_39, %c_ptrs_38 : tensor<32x1x!tt.ptr, #blocked>, tensor<32x1xi32, #blocked> loc(#loc94)\n %c_ptrs_41 = tt.expand_dims %offs_bn_15 {axis = 0 : i32} : tensor<16xi32, #ttg.slice<{dim = 0, parent = #blocked}>> -> tensor<1x16xi32, #blocked> loc(#loc95)\n %c_ptrs_42 = tt.broadcast %c_ptrs_40 : tensor<32x1x!tt.ptr, #blocked> -> tensor<32x16x!tt.ptr, #blocked> loc(#loc94)\n %c_ptrs_43 = tt.broadcast %c_ptrs_41 : tensor<1x16xi32, #blocked> -> tensor<32x16xi32, #blocked> loc(#loc94)\n %c_ptrs_44 = tt.addptr %c_ptrs_42, %c_ptrs_43 : tensor<32x16x!tt.ptr, #blocked>, tensor<32x16xi32, #blocked> loc(#loc94)\n %c_mask = tt.splat %M : i32 -> tensor<32x1xi32, #blocked> loc(#loc96)\n %c_mask_45 = arith.cmpi slt, %c_ptrs, %c_mask : tensor<32x1xi32, #blocked> loc(#loc96)\n %c_mask_46 = tt.splat %N : i32 -> tensor<1x16xi32, #blocked> loc(#loc97)\n %c_mask_47 = arith.cmpi slt, %c_ptrs_41, %c_mask_46 : tensor<1x16xi32, #blocked> loc(#loc97)\n %c_mask_48 = tt.broadcast %c_mask_45 : tensor<32x1xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc98)\n %c_mask_49 = tt.broadcast %c_mask_47 : tensor<1x16xi1, #blocked> -> tensor<32x16xi1, #blocked> loc(#loc98)\n %c_mask_50 = arith.andi %c_mask_48, %c_mask_49 : tensor<32x16xi1, #blocked> loc(#loc98)\n %2 = ttg.convert_layout %c_block : tensor<32x16xf16, #mma> -> tensor<32x16xf16, #blocked> loc(#loc46)\n tt.store %c_ptrs_44, %2, %c_mask_50 : tensor<32x16x!tt.ptr, #blocked> loc(#loc46)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":311:15)\n#loc3 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:13)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":312:21)\n#loc5 = loc(\"/python/lib/python3.14/site-packages/triton/language/standard.py\":43:12)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":313:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":315:20)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:26)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":317:22)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:32)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":318:17)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:18)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":319:17)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:20)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:43)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":321:19)\n#loc17 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:20)\n#loc18 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:43)\n#loc19 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":322:19)\n#loc20 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:23)\n#loc21 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:54)\n#loc22 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":324:18)\n#loc23 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:23)\n#loc24 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:53)\n#loc25 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":325:18)\n#loc26 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:27)\n#loc27 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:23)\n#loc28 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":337:13)\n#loc29 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":328:9)\n#loc30 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:52)\n#loc31 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:48)\n#loc32 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":330:30)\n#loc33 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":329:23)\n#loc34 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":333:30)\n#loc35 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":332:23)\n#loc36 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":335:28)\n#loc37 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":336:13)\n#loc38 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":338:19)\n#loc39 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:34)\n#loc40 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:22)\n#loc41 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:18)\n#loc42 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":342:65)\n#loc43 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:19)\n#loc44 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:44)\n#loc45 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":343:18)\n#loc46 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":344:9)\n#loc56 = loc(\"pid\"(#loc2))\n#loc57 = loc(\"num_pid_m\"(#loc4))\n#loc58 = loc(\"num_pid_n\"(#loc6))\n#loc59 = loc(\"group_id\"(#loc7))\n#loc60 = loc(\"group_size\"(#loc8))\n#loc61 = loc(\"group_size\"(#loc9))\n#loc62 = loc(\"pid_m\"(#loc10))\n#loc63 = loc(\"pid_m\"(#loc11))\n#loc64 = loc(\"pid_n\"(#loc12))\n#loc65 = loc(\"pid_n\"(#loc13))\n#loc66 = loc(\"offs_am\"(#loc14))\n#loc67 = loc(\"offs_am\"(#loc15))\n#loc68 = loc(\"offs_am\"(#loc16))\n#loc69 = loc(\"offs_bn\"(#loc17))\n#loc70 = loc(\"offs_bn\"(#loc18))\n#loc71 = loc(\"offs_bn\"(#loc19))\n#loc72 = loc(\"a_ptrs\"(#loc20))\n#loc73 = loc(\"a_ptrs\"(#loc21))\n#loc74 = loc(\"a_ptrs\"(#loc22))\n#loc75 = loc(\"b_ptrs\"(#loc23))\n#loc76 = loc(\"b_ptrs\"(#loc24))\n#loc77 = loc(\"b_ptrs\"(#loc25))\n#loc78 = loc(callsite(#loc3 at #loc26))\n#loc79 = loc(callsite(#loc5 at #loc26))\n#loc80 = loc(\"b_ptrs\"(#loc27))\n#loc81 = loc(\"b_ptrs\"(#loc28))\n#loc82 = loc(\"a_ptrs\"(#loc29))\n#loc83 = loc(\"a_block\"(#loc30))\n#loc84 = loc(\"a_block\"(#loc31))\n#loc85 = loc(\"a_block\"(#loc32))\n#loc86 = loc(\"a_block\"(#loc33))\n#loc87 = loc(\"b_block\"(#loc34))\n#loc88 = loc(\"b_block\"(#loc35))\n#loc89 = loc(\"accumulator\"(#loc36))\n#loc90 = loc(\"a_ptrs\"(#loc37))\n#loc91 = loc(\"c_block\"(#loc38))\n#loc92 = loc(\"c_ptrs\"(#loc39))\n#loc93 = loc(\"c_ptrs\"(#loc40))\n#loc94 = loc(\"c_ptrs\"(#loc41))\n#loc95 = loc(\"c_ptrs\"(#loc42))\n#loc96 = loc(\"c_mask\"(#loc43))\n#loc97 = loc(\"c_mask\"(#loc44))\n#loc98 = loc(\"c_mask\"(#loc45))\n#loc99 = loc(callsite(#loc3 at #loc57))\n#loc100 = loc(callsite(#loc5 at #loc57))\n#loc101 = loc(callsite(#loc3 at #loc58))\n#loc102 = loc(callsite(#loc5 at #loc58))\n#loc103 = loc(\"b_ptrs\"(#loc82))\n#loc104 = loc(\"accumulator\"(#loc103))\n","matmul_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n@global_smem = external addrspace(3) global [0 x i8], align 16\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @matmul_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i32 %8, ptr addrspace(1) readnone captures(none) %9, ptr addrspace(1) readnone captures(none) %10) local_unnamed_addr #0 !dbg !4 {\n %12 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %13 = add i32 %3, 31, !dbg !13\n %14 = sdiv i32 %13, 32, !dbg !18\n %15 = add i32 %4, 15, !dbg !19\n %16 = sdiv i32 %15, 16, !dbg !21\n %.frozen = freeze i32 %16, !dbg !22\n %17 = sdiv i32 %12, %.frozen, !dbg !22\n %18 = sub i32 %14, %17, !dbg !23\n %19 = tail call i32 @llvm.smin.i32(i32 %18, i32 1), !dbg !24\n %20 = srem i32 %12, %19, !dbg !25\n %21 = add i32 %20, %17, !dbg !26\n %22 = mul i32 %17, %.frozen, !dbg !27\n %.decomposed = sub i32 %12, %22, !dbg !27\n %23 = sdiv i32 %.decomposed, %19, !dbg !28\n %24 = shl i32 %21, 5, !dbg !29\n %25 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !30\n %26 = lshr i32 %25, 1, !dbg !30\n %27 = and i32 %26, 15, !dbg !30\n %28 = or disjoint i32 %24, %27, !dbg !29\n %29 = or disjoint i32 %28, 16, !dbg !29\n %30 = shl nsw i32 %23, 4, !dbg !31\n %31 = and i32 %25, 1, !dbg !32\n %32 = shl nuw nsw i32 %31, 3, !dbg !32\n %33 = or disjoint i32 %30, %32, !dbg !31\n %34 = add i32 %5, 15, !dbg !33\n %35 = sdiv i32 %34, 16, !dbg !35\n %36 = icmp sgt i32 %34, 15, !dbg !36\n br i1 %36, label %.lr.ph, label %.._crit_edge_crit_edge, !dbg !36\n\n.._crit_edge_crit_edge: ; preds = %11\n %.pre = shl nuw nsw i32 %25, 4, !dbg !37\n %.pre109 = and i32 %25, 4, !dbg !37\n %.pre111 = and i32 %25, 8, !dbg !37\n br label %._crit_edge, !dbg !36\n\n.lr.ph: ; preds = %11\n %37 = shl i32 %7, 4, !dbg !38\n %38 = srem i32 %33, %4, !dbg !39\n %39 = mul i32 %7, %27, !dbg !40\n %40 = add i32 %38, %39, !dbg !40\n %41 = sext i32 %40 to i64, !dbg !41\n %42 = getelementptr [2 x i8], ptr addrspace(1) %1, i64 %41, !dbg !41\n %43 = srem i32 %29, %3, !dbg !42\n %44 = mul i32 %43, %6, !dbg !43\n %45 = add i32 %44, %32, !dbg !43\n %46 = sext i32 %45 to i64, !dbg !44\n %47 = getelementptr [2 x i8], ptr addrspace(1) %0, i64 %46, !dbg !44\n %48 = srem i32 %28, %3, !dbg !42\n %49 = mul i32 %48, %6, !dbg !43\n %50 = add i32 %49, %32, !dbg !43\n %51 = sext i32 %50 to i64, !dbg !44\n %52 = getelementptr [2 x i8], ptr addrspace(1) %0, i64 %51, !dbg !44\n %53 = shl nuw nsw i32 %25, 4\n %54 = and i32 %53, 368\n %55 = and i32 %25, 8\n %56 = icmp eq i32 %55, 0\n %57 = select i1 %56, i32 0, i32 144\n %58 = xor i32 %57, %54\n %59 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %58\n %60 = getelementptr inbounds nuw i8, ptr addrspace(3) %59, i32 512\n %61 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 1024), i32 %58\n %62 = shl nuw nsw i32 %25, 5\n %63 = and i32 %62, 352\n %64 = and i32 %25, 4\n %65 = icmp eq i32 %64, 0\n %66 = select i1 %65, i32 0, i32 144\n %67 = and i32 %25, 16\n %68 = or disjoint i32 %66, %63\n %69 = xor i32 %68, %67\n %70 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %69\n %71 = getelementptr inbounds nuw i8, ptr addrspace(3) %70, i32 512\n %72 = xor i32 %66, %67\n %73 = getelementptr inbounds nuw i8, ptr addrspace(3) getelementptr (i8, ptr addrspace(3) @global_smem, i32 1024), i32 %63\n %74 = getelementptr inbounds nuw i8, ptr addrspace(3) %73, i32 %72\n %75 = sext i32 %37 to i64\n br label %76, !dbg !36\n\n76: ; preds = %.lr.ph, %76\n %.pn47107 = phi ptr addrspace(1) [ %42, %.lr.ph ], [ %150, %76 ]\n %.pn15106 = phi ptr addrspace(1) [ %47, %.lr.ph ], [ %149, %76 ]\n %.pn31105 = phi ptr addrspace(1) [ %52, %.lr.ph ], [ %148, %76 ]\n %77 = phi float [ 0.000000e+00, %.lr.ph ], [ %129, %76 ]\n %78 = phi float [ 0.000000e+00, %.lr.ph ], [ %130, %76 ]\n %79 = phi float [ 0.000000e+00, %.lr.ph ], [ %131, %76 ]\n %80 = phi float [ 0.000000e+00, %.lr.ph ], [ %132, %76 ]\n %81 = phi float [ 0.000000e+00, %.lr.ph ], [ %134, %76 ]\n %82 = phi float [ 0.000000e+00, %.lr.ph ], [ %135, %76 ]\n %83 = phi float [ 0.000000e+00, %.lr.ph ], [ %136, %76 ]\n %84 = phi float [ 0.000000e+00, %.lr.ph ], [ %137, %76 ]\n %85 = phi float [ 0.000000e+00, %.lr.ph ], [ %139, %76 ]\n %86 = phi float [ 0.000000e+00, %.lr.ph ], [ %140, %76 ]\n %87 = phi float [ 0.000000e+00, %.lr.ph ], [ %141, %76 ]\n %88 = phi float [ 0.000000e+00, %.lr.ph ], [ %142, %76 ]\n %89 = phi float [ 0.000000e+00, %.lr.ph ], [ %144, %76 ]\n %90 = phi float [ 0.000000e+00, %.lr.ph ], [ %145, %76 ]\n %91 = phi float [ 0.000000e+00, %.lr.ph ], [ %146, %76 ]\n %92 = phi float [ 0.000000e+00, %.lr.ph ], [ %147, %76 ]\n %93 = phi i32 [ 0, %.lr.ph ], [ %151, %76 ]\n %94 = shl i32 %93, 4, !dbg !45\n %95 = sub i32 %5, %94, !dbg !46\n %96 = icmp slt i32 %32, %95, !dbg !47\n %97 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn31105, i1 %96) #6, !dbg !48\n %98 = extractvalue { i32, i32, i32, i32 } %97, 0, !dbg !48\n %99 = extractvalue { i32, i32, i32, i32 } %97, 1, !dbg !48\n %100 = extractvalue { i32, i32, i32, i32 } %97, 2, !dbg !48\n %101 = extractvalue { i32, i32, i32, i32 } %97, 3, !dbg !48\n %102 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn15106, i1 %96) #6, !dbg !48\n %103 = extractvalue { i32, i32, i32, i32 } %102, 0, !dbg !48\n %104 = extractvalue { i32, i32, i32, i32 } %102, 1, !dbg !48\n %105 = extractvalue { i32, i32, i32, i32 } %102, 2, !dbg !48\n %106 = extractvalue { i32, i32, i32, i32 } %102, 3, !dbg !48\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !48\n tail call void asm sideeffect \"st.shared::cta.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %59, i32 %98, i32 %99, i32 %100, i32 %101) #6, !dbg !48\n tail call void asm sideeffect \"st.shared::cta.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) nonnull %60, i32 %103, i32 %104, i32 %105, i32 %106) #6, !dbg !48\n %107 = icmp slt i32 %27, %95, !dbg !49\n %108 = tail call { i32, i32, i32, i32 } asm sideeffect \"mov.u32 $0, $4;\\0A\\09mov.u32 $1, $5;\\0A\\09mov.u32 $2, $6;\\0A\\09mov.u32 $3, $7;\\0A\\09@$9 ld.global.v4.b32 { $0, $1, $2, $3 }, [ $8 + 0 ];\", \"=r,=r,=r,=r,r,r,r,r,l,b\"(i32 0, i32 0, i32 0, i32 0, ptr addrspace(1) %.pn47107, i1 %107) #6, !dbg !50\n %109 = extractvalue { i32, i32, i32, i32 } %108, 0, !dbg !50\n %110 = extractvalue { i32, i32, i32, i32 } %108, 1, !dbg !50\n %111 = extractvalue { i32, i32, i32, i32 } %108, 2, !dbg !50\n %112 = extractvalue { i32, i32, i32, i32 } %108, 3, !dbg !50\n tail call void asm sideeffect \"st.shared::cta.v4.b32 [ $0 + 0 ], { $1, $2, $3, $4 };\", \"r,r,r,r,r\"(ptr addrspace(3) %61, i32 %109, i32 %110, i32 %111, i32 %112) #6, !dbg !50\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !48\n %113 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %70), !dbg !48\n %114 = extractvalue { i32, i32, i32, i32 } %113, 0, !dbg !48\n %115 = extractvalue { i32, i32, i32, i32 } %113, 1, !dbg !48\n %116 = extractvalue { i32, i32, i32, i32 } %113, 2, !dbg !48\n %117 = extractvalue { i32, i32, i32, i32 } %113, 3, !dbg !48\n %118 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) nonnull %71), !dbg !48\n %119 = extractvalue { i32, i32, i32, i32 } %118, 0, !dbg !48\n %120 = extractvalue { i32, i32, i32, i32 } %118, 1, !dbg !48\n %121 = extractvalue { i32, i32, i32, i32 } %118, 2, !dbg !48\n %122 = extractvalue { i32, i32, i32, i32 } %118, 3, !dbg !48\n %123 = tail call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %74), !dbg !50\n %124 = extractvalue { i32, i32, i32, i32 } %123, 0, !dbg !50\n %125 = extractvalue { i32, i32, i32, i32 } %123, 1, !dbg !50\n %126 = extractvalue { i32, i32, i32, i32 } %123, 2, !dbg !50\n %127 = extractvalue { i32, i32, i32, i32 } %123, 3, !dbg !50\n %128 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %77, float %78, float %79, float %80, i32 %114, i32 %115, i32 %116, i32 %117, i32 %124, i32 %125) #6, !dbg !51\n %129 = extractvalue { float, float, float, float } %128, 0, !dbg !51\n %130 = extractvalue { float, float, float, float } %128, 1, !dbg !51\n %131 = extractvalue { float, float, float, float } %128, 2, !dbg !51\n %132 = extractvalue { float, float, float, float } %128, 3, !dbg !51\n %133 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %81, float %82, float %83, float %84, i32 %114, i32 %115, i32 %116, i32 %117, i32 %126, i32 %127) #6, !dbg !51\n %134 = extractvalue { float, float, float, float } %133, 0, !dbg !51\n %135 = extractvalue { float, float, float, float } %133, 1, !dbg !51\n %136 = extractvalue { float, float, float, float } %133, 2, !dbg !51\n %137 = extractvalue { float, float, float, float } %133, 3, !dbg !51\n %138 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %85, float %86, float %87, float %88, i32 %119, i32 %120, i32 %121, i32 %122, i32 %124, i32 %125) #6, !dbg !51\n %139 = extractvalue { float, float, float, float } %138, 0, !dbg !51\n %140 = extractvalue { float, float, float, float } %138, 1, !dbg !51\n %141 = extractvalue { float, float, float, float } %138, 2, !dbg !51\n %142 = extractvalue { float, float, float, float } %138, 3, !dbg !51\n %143 = tail call { float, float, float, float } asm sideeffect \"mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { $0, $1, $2, $3 }, { $8, $9, $10, $11 }, { $12, $13 }, { $4, $5, $6, $7 };\", \"=f,=f,=f,=f,0,1,2,3,r,r,r,r,r,r\"(float %89, float %90, float %91, float %92, i32 %119, i32 %120, i32 %121, i32 %122, i32 %126, i32 %127) #6, !dbg !51\n %144 = extractvalue { float, float, float, float } %143, 0, !dbg !51\n %145 = extractvalue { float, float, float, float } %143, 1, !dbg !51\n %146 = extractvalue { float, float, float, float } %143, 2, !dbg !51\n %147 = extractvalue { float, float, float, float } %143, 3, !dbg !51\n %148 = getelementptr i8, ptr addrspace(1) %.pn31105, i64 32, !dbg !52\n %149 = getelementptr i8, ptr addrspace(1) %.pn15106, i64 32, !dbg !52\n %150 = getelementptr [2 x i8], ptr addrspace(1) %.pn47107, i64 %75, !dbg !53\n %151 = add nuw nsw i32 %93, 1, !dbg !36\n %exitcond.not = icmp eq i32 %151, %35, !dbg !36\n br i1 %exitcond.not, label %._crit_edge.loopexit, label %76, !dbg !36\n\n._crit_edge.loopexit: ; preds = %76\n %152 = insertelement <16 x float> poison, float %129, i64 0, !dbg !54\n %153 = insertelement <16 x float> %152, float %130, i64 1, !dbg !54\n %154 = insertelement <16 x float> %153, float %131, i64 2, !dbg !54\n %155 = insertelement <16 x float> %154, float %132, i64 3, !dbg !54\n %156 = insertelement <16 x float> %155, float %134, i64 4, !dbg !54\n %157 = insertelement <16 x float> %156, float %135, i64 5, !dbg !54\n %158 = insertelement <16 x float> %157, float %136, i64 6, !dbg !54\n %159 = insertelement <16 x float> %158, float %137, i64 7, !dbg !54\n %160 = insertelement <16 x float> %159, float %139, i64 8, !dbg !54\n %161 = insertelement <16 x float> %160, float %140, i64 9, !dbg !54\n %162 = insertelement <16 x float> %161, float %141, i64 10, !dbg !54\n %163 = insertelement <16 x float> %162, float %142, i64 11, !dbg !54\n %164 = insertelement <16 x float> %163, float %144, i64 12, !dbg !54\n %165 = insertelement <16 x float> %164, float %145, i64 13, !dbg !54\n %166 = insertelement <16 x float> %165, float %146, i64 14, !dbg !54\n %167 = insertelement <16 x float> %166, float %147, i64 15, !dbg !54\n %168 = fptrunc <16 x float> %167 to <16 x half>, !dbg !54\n %169 = bitcast <16 x half> %168 to <8 x i32>, !dbg !37\n br label %._crit_edge, !dbg !54\n\n._crit_edge: ; preds = %._crit_edge.loopexit, %.._crit_edge_crit_edge\n %.pre-phi112 = phi i32 [ %.pre111, %.._crit_edge_crit_edge ], [ %55, %._crit_edge.loopexit ], !dbg !37\n %.pre-phi110 = phi i32 [ %.pre109, %.._crit_edge_crit_edge ], [ %64, %._crit_edge.loopexit ], !dbg !37\n %.pre-phi = phi i32 [ %.pre, %.._crit_edge_crit_edge ], [ %53, %._crit_edge.loopexit ], !dbg !37\n %.bc142 = phi <8 x i32> [ zeroinitializer, %.._crit_edge_crit_edge ], [ %169, %._crit_edge.loopexit ], !dbg !55\n %170 = mul i32 %28, %8, !dbg !56\n %171 = mul i32 %8, %29, !dbg !56\n %172 = sext i32 %170 to i64, !dbg !57\n %173 = getelementptr [2 x i8], ptr addrspace(1) %2, i64 %172, !dbg !57\n %174 = sext i32 %171 to i64, !dbg !57\n %175 = getelementptr [2 x i8], ptr addrspace(1) %2, i64 %174, !dbg !57\n %176 = sext i32 %33 to i64, !dbg !57\n %177 = getelementptr [2 x i8], ptr addrspace(1) %173, i64 %176, !dbg !57\n %178 = getelementptr [2 x i8], ptr addrspace(1) %175, i64 %176, !dbg !57\n %179 = icmp slt i32 %28, %3, !dbg !58\n %180 = icmp slt i32 %29, %3, !dbg !58\n %181 = icmp slt i32 %33, %4, !dbg !59\n %182 = and i1 %179, %181, !dbg !60\n %183 = and i1 %180, %181, !dbg !60\n tail call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0), !dbg !37\n %184 = and i32 %.pre-phi, 176, !dbg !37\n %185 = icmp eq i32 %.pre-phi110, 0, !dbg !37\n %186 = select i1 %185, i32 0, i32 320, !dbg !37\n %187 = shl nuw nsw i32 %25, 2, !dbg !37\n %188 = and i32 %187, 64, !dbg !37\n %189 = or disjoint i32 %186, %184, !dbg !37\n %190 = xor i32 %189, %188, !dbg !37\n %191 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %190, !dbg !37\n %.extract129 = extractelement <8 x i32> %.bc142, i64 0, !dbg !37\n %.extract131 = extractelement <8 x i32> %.bc142, i64 1, !dbg !37\n %.extract133 = extractelement <8 x i32> %.bc142, i64 2, !dbg !37\n %.extract135 = extractelement <8 x i32> %.bc142, i64 3, !dbg !37\n tail call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %191, i32 %.extract129, i32 %.extract131, i32 %.extract133, i32 %.extract135), !dbg !37\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !37\n %192 = shl nuw nsw i32 %25, 3, !dbg !37\n %193 = and i32 %192, 176, !dbg !37\n %194 = shl nuw nsw i32 %31, 6, !dbg !37\n %195 = icmp eq i32 %.pre-phi112, 0, !dbg !37\n %196 = select i1 %195, i32 0, i32 320, !dbg !37\n %197 = or disjoint i32 %193, %194, !dbg !37\n %198 = xor i32 %197, %196, !dbg !37\n %199 = getelementptr inbounds nuw i8, ptr addrspace(3) @global_smem, i32 %198, !dbg !37\n %200 = load <4 x i32>, ptr addrspace(3) %199, align 16, !dbg !37\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !37\n %.extract137 = extractelement <8 x i32> %.bc142, i64 4, !dbg !37\n %.extract139 = extractelement <8 x i32> %.bc142, i64 5, !dbg !37\n %.extract141 = extractelement <8 x i32> %.bc142, i64 6, !dbg !37\n %.extract143 = extractelement <8 x i32> %.bc142, i64 7, !dbg !37\n tail call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %191, i32 %.extract137, i32 %.extract139, i32 %.extract141, i32 %.extract143), !dbg !37\n tail call void @llvm.nvvm.bar.warp.sync(i32 -1), !dbg !37\n %201 = load <4 x i32>, ptr addrspace(3) %199, align 16, !dbg !37\n %.extract = extractelement <4 x i32> %200, i64 0, !dbg !37\n %.extract92 = extractelement <4 x i32> %200, i64 1, !dbg !37\n %.extract94 = extractelement <4 x i32> %200, i64 2, !dbg !37\n %.extract96 = extractelement <4 x i32> %200, i64 3, !dbg !37\n tail call void asm sideeffect \"@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };\", \"r,r,r,r,l,b\"(i32 %.extract, i32 %.extract92, i32 %.extract94, i32 %.extract96, ptr addrspace(1) %177, i1 %182) #6, !dbg !37\n %.extract98 = extractelement <4 x i32> %201, i64 0, !dbg !37\n %.extract100 = extractelement <4 x i32> %201, i64 1, !dbg !37\n %.extract102 = extractelement <4 x i32> %201, i64 2, !dbg !37\n %.extract104 = extractelement <4 x i32> %201, i64 3, !dbg !37\n tail call void asm sideeffect \"@$5 st.global.v4.b32 [ $4 + 0 ], { $0, $1, $2, $3 };\", \"r,r,r,r,l,b\"(i32 %.extract98, i32 %.extract100, i32 %.extract102, i32 %.extract104, ptr addrspace(1) %178, i1 %183) #6, !dbg !37\n ret void, !dbg !61\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none)\ndeclare i32 @llvm.smin.i32(i32, i32) #2\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.barrier.cta.sync.aligned.all(i32) #3\n\n; Function Attrs: convergent nocallback nounwind memory(argmem: write)\ndeclare void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) writeonly captures(none), i32, i32, i32, i32) #4\n\n; Function Attrs: convergent nocallback nounwind\ndeclare void @llvm.nvvm.bar.warp.sync(i32) #3\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\n; Function Attrs: convergent nocallback nofree nounwind memory(argmem: read)\ndeclare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) readonly captures(none)) #5\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"32\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { mustprogress nocallback nocreateundeforpoison nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #3 = { convergent nocallback nounwind }\nattributes #4 = { convergent nocallback nounwind memory(argmem: write) }\nattributes #5 = { convergent nocallback nofree nounwind memory(argmem: read) }\nattributes #6 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"matmul_kernel\", linkageName: \"matmul_kernel\", scope: !1, file: !1, line: 292, type: !5, scopeLine: 292, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !9, !9, !9, !9, !9, !9, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"half\", size: 16, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 311, column: 15, scope: !4)\n!13 = !DILocation(line: 43, column: 13, scope: !14, inlinedAt: !16)\n!14 = distinct !DILexicalBlockFile(scope: !4, file: !15, discriminator: 0)\n!15 = !DIFile(filename: \"standard.py\", directory: \"/python/lib/python3.14/site-packages/triton/language\")\n!16 = !DILocation(line: 312, column: 21, scope: !17)\n!17 = distinct !DILexicalBlockFile(scope: !4, file: !1, discriminator: 0)\n!18 = !DILocation(line: 43, column: 12, scope: !14, inlinedAt: !16)\n!19 = !DILocation(line: 43, column: 13, scope: !14, inlinedAt: !20)\n!20 = !DILocation(line: 313, column: 21, scope: !17)\n!21 = !DILocation(line: 43, column: 12, scope: !14, inlinedAt: !20)\n!22 = !DILocation(line: 315, column: 20, scope: !4)\n!23 = !DILocation(line: 317, column: 26, scope: !4)\n!24 = !DILocation(line: 317, column: 22, scope: !4)\n!25 = !DILocation(line: 318, column: 32, scope: !4)\n!26 = !DILocation(line: 318, column: 17, scope: !4)\n!27 = !DILocation(line: 319, column: 18, scope: !4)\n!28 = !DILocation(line: 319, column: 17, scope: !4)\n!29 = !DILocation(line: 321, column: 20, scope: !4)\n!30 = !DILocation(line: 321, column: 43, scope: !4)\n!31 = !DILocation(line: 322, column: 20, scope: !4)\n!32 = !DILocation(line: 322, column: 43, scope: !4)\n!33 = !DILocation(line: 43, column: 13, scope: !14, inlinedAt: !34)\n!34 = !DILocation(line: 328, column: 27, scope: !17)\n!35 = !DILocation(line: 43, column: 12, scope: !14, inlinedAt: !34)\n!36 = !DILocation(line: 328, column: 9, scope: !4)\n!37 = !DILocation(line: 344, column: 9, scope: !4)\n!38 = !DILocation(line: 337, column: 23, scope: !4)\n!39 = !DILocation(line: 322, column: 19, scope: !4)\n!40 = !DILocation(line: 325, column: 23, scope: !4)\n!41 = !DILocation(line: 325, column: 18, scope: !4)\n!42 = !DILocation(line: 321, column: 19, scope: !4)\n!43 = !DILocation(line: 324, column: 23, scope: !4)\n!44 = !DILocation(line: 324, column: 18, scope: !4)\n!45 = !DILocation(line: 330, column: 52, scope: !4)\n!46 = !DILocation(line: 330, column: 48, scope: !4)\n!47 = !DILocation(line: 330, column: 30, scope: !4)\n!48 = !DILocation(line: 329, column: 23, scope: !4)\n!49 = !DILocation(line: 333, column: 30, scope: !4)\n!50 = !DILocation(line: 332, column: 23, scope: !4)\n!51 = !DILocation(line: 335, column: 28, scope: !4)\n!52 = !DILocation(line: 336, column: 13, scope: !4)\n!53 = !DILocation(line: 337, column: 13, scope: !4)\n!54 = !DILocation(line: 338, column: 19, scope: !4)\n!55 = !DILocation(line: 327, column: 23, scope: !4)\n!56 = !DILocation(line: 342, column: 22, scope: !4)\n!57 = !DILocation(line: 342, column: 18, scope: !4)\n!58 = !DILocation(line: 343, column: 19, scope: !4)\n!59 = !DILocation(line: 343, column: 44, scope: !4)\n!60 = !DILocation(line: 343, column: 18, scope: !4)\n!61 = !DILocation(line: 292, column: 5, scope: !4)\n","matmul_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.8\n.target sm_90a\n.address_size 64\n\n\t// .globl\tmatmul_kernel // -- Begin function matmul_kernel\n.extern .shared .align 16 .b8 global_smem[];\n // @matmul_kernel\n.visible .entry matmul_kernel(\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_2,\n\t.param .u32 matmul_kernel_param_3,\n\t.param .u32 matmul_kernel_param_4,\n\t.param .u32 matmul_kernel_param_5,\n\t.param .u32 matmul_kernel_param_6,\n\t.param .u32 matmul_kernel_param_7,\n\t.param .u32 matmul_kernel_param_8,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_9,\n\t.param .u64 .ptr .global .align 1 matmul_kernel_param_10\n)\n.reqntid 32\n{\n\t.reg .pred \t%p<12>;\n\t.reg .b32 \t%r<144>;\n\t.reg .b64 \t%rd<13>;\n\t.loc\t1 292 0 // generate_examples.py:292:0\n$L__func_begin0:\n\t.loc\t1 292 0 // generate_examples.py:292:0\n\n// %bb.0:\n\tld.param.b32 \t%r15, [matmul_kernel_param_8];\n\tld.param.b32 \t%r115, [matmul_kernel_param_5];\n\tld.param.b32 \t%r12, [matmul_kernel_param_4];\n\tld.param.b32 \t%r11, [matmul_kernel_param_3];\n\tld.param.b64 \t%rd4, [matmul_kernel_param_2];\n$L__tmp0:\n\t.loc\t1 311 15 // generate_examples.py:311:15\n\tmov.u32 \t%r16, %ctaid.x;\n$L__tmp1:\n\t.loc\t1 312 21 // generate_examples.py:312:21\n\t.loc\t2 43 13, function_name $L__info_string0, inlined_at 1 312 21 // standard.py:43:13\n\tadd.s32 \t%r17, %r11, 31;\n\t.loc\t2 43 12, function_name $L__info_string0, inlined_at 1 312 21 // standard.py:43:12\n\tshr.s32 \t%r18, %r17, 31;\n\tshr.u32 \t%r19, %r18, 27;\n\tadd.s32 \t%r20, %r17, %r19;\n\tshr.s32 \t%r21, %r20, 5;\n$L__tmp2:\n\t.loc\t1 313 21 // generate_examples.py:313:21\n\t.loc\t2 43 13, function_name $L__info_string0, inlined_at 1 313 21 // standard.py:43:13\n\tadd.s32 \t%r22, %r12, 15;\n\t.loc\t2 43 12, function_name $L__info_string0, inlined_at 1 313 21 // standard.py:43:12\n\tshr.s32 \t%r23, %r22, 31;\n\tshr.u32 \t%r24, %r23, 28;\n\tadd.s32 \t%r25, %r22, %r24;\n\tshr.s32 \t%r26, %r25, 4;\n$L__tmp3:\n\t.loc\t1 315 20 // generate_examples.py:315:20\n\tdiv.s32 \t%r27, %r16, %r26;\n\t.loc\t1 317 26 // generate_examples.py:317:26\n\tsub.s32 \t%r28, %r21, %r27;\n\t.loc\t1 317 22 // generate_examples.py:317:22\n\tmin.s32 \t%r29, %r28, 1;\n\t.loc\t1 318 32 // generate_examples.py:318:32\n\trem.s32 \t%r30, %r16, %r29;\n\t.loc\t1 318 17 // generate_examples.py:318:17\n\tadd.s32 \t%r31, %r30, %r27;\n\t.loc\t1 319 18 // generate_examples.py:319:18\n\tmul.lo.s32 \t%r32, %r27, %r26;\n\tsub.s32 \t%r33, %r16, %r32;\n\t.loc\t1 319 17 // generate_examples.py:319:17\n\tdiv.s32 \t%r34, %r33, %r29;\n\t.loc\t1 321 20 // generate_examples.py:321:20\n\tshl.b32 \t%r35, %r31, 5;\n\t.loc\t1 321 43 // generate_examples.py:321:43\n\tmov.u32 \t%r1, %tid.x;\n\tbfe.u32 \t%r2, %r1, 1, 4;\n\t.loc\t1 321 20 // generate_examples.py:321:20\n\tor.b32 \t%r3, %r35, %r2;\n\tor.b32 \t%r4, %r3, 16;\n\t.loc\t1 322 20 // generate_examples.py:322:20\n\tshl.b32 \t%r36, %r34, 4;\n\t.loc\t1 322 43 // generate_examples.py:322:43\n\tand.b32 \t%r5, %r1, 1;\n\tshl.b32 \t%r6, %r5, 3;\n\t.loc\t1 322 20 // generate_examples.py:322:20\n\tor.b32 \t%r7, %r36, %r6;\n$L__tmp4:\n\t.loc\t1 328 27 // generate_examples.py:328:27\n\t.loc\t2 43 13, function_name $L__info_string0, inlined_at 1 328 27 // standard.py:43:13\n\tadd.s32 \t%r37, %r115, 15;\n$L__tmp5:\n\t.loc\t1 328 9 // generate_examples.py:328:9\n\tsetp.gt.s32 \t%p1, %r37, 15;\n\tmov.b32 \t%r114, global_smem;\n\t@!%p1 bra \t$L__BB0_1;\n// %bb.2: // %.lr.ph\n\t.loc\t1 0 9 // generate_examples.py:0:9\n\tld.param.b32 \t%r14, [matmul_kernel_param_7];\n\tld.param.b32 \t%r13, [matmul_kernel_param_6];\n\tld.param.b64 \t%rd3, [matmul_kernel_param_1];\n\tld.param.b64 \t%rd2, [matmul_kernel_param_0];\n\tshr.s32 \t%r38, %r37, 31;\n\tshr.u32 \t%r39, %r38, 28;\n\tadd.s32 \t%r40, %r37, %r39;\n\tshr.s32 \t%r116, %r40, 4;\n\t.loc\t1 337 23 // generate_examples.py:337:23\n\tshl.b32 \t%r41, %r14, 4;\n\t.loc\t1 322 19 // generate_examples.py:322:19\n\trem.s32 \t%r42, %r7, %r12;\n\t.loc\t1 325 23 // generate_examples.py:325:23\n\tmad.lo.s32 \t%r43, %r14, %r2, %r42;\n\t.loc\t1 325 18 // generate_examples.py:325:18\n\tmad.wide.s32 \t%rd10, %r43, 2, %rd3;\n\t.loc\t1 321 19 // generate_examples.py:321:19\n\trem.s32 \t%r44, %r4, %r11;\n\t.loc\t1 324 23 // generate_examples.py:324:23\n\tmad.lo.s32 \t%r45, %r44, %r13, %r6;\n\t.loc\t1 324 18 // generate_examples.py:324:18\n\tmad.wide.s32 \t%rd11, %r45, 2, %rd2;\n\t.loc\t1 321 19 // generate_examples.py:321:19\n\trem.s32 \t%r46, %r3, %r11;\n\t.loc\t1 324 23 // generate_examples.py:324:23\n\tmad.lo.s32 \t%r47, %r46, %r13, %r6;\n\t.loc\t1 324 18 // generate_examples.py:324:18\n\tmad.wide.s32 \t%rd12, %r47, 2, %rd2;\n\tshl.b32 \t%r135, %r1, 4;\n\tand.b32 \t%r48, %r135, 368;\n\tand.b32 \t%r133, %r1, 8;\n\tbfe.s32 \t%r49, %r1, 3, 1;\n\tand.b32 \t%r50, %r49, 144;\n\txor.b32 \t%r51, %r50, %r48;\n\tadd.s32 \t%r71, %r114, %r51;\n\tadd.s32 \t%r72, %r71, 512;\n\tadd.s32 \t%r77, %r71, 1024;\n\tshl.b32 \t%r52, %r1, 5;\n\tand.b32 \t%r53, %r52, 352;\n\tand.b32 \t%r134, %r1, 4;\n\tbfe.s32 \t%r54, %r1, 2, 1;\n\tand.b32 \t%r55, %r54, 144;\n\tand.b32 \t%r56, %r1, 16;\n\tor.b32 \t%r57, %r55, %r53;\n\txor.b32 \t%r58, %r57, %r56;\n\tadd.s32 \t%r8, %r114, %r58;\n\tadd.s32 \t%r9, %r8, 512;\n\txor.b32 \t%r59, %r55, %r56;\n\tadd.s32 \t%r60, %r114, %r53;\n\tadd.s32 \t%r61, %r60, %r59;\n\tadd.s32 \t%r10, %r61, 1024;\n\t.loc\t1 328 9 // generate_examples.py:328:9\n\tmul.wide.s32 \t%rd1, %r41, 2;\n\tmov.b32 \t%r117, 0f00000000;\n\tmov.b32 \t%r118, %r117;\n\tmov.b32 \t%r119, %r117;\n\tmov.b32 \t%r120, %r117;\n\tmov.b32 \t%r121, %r117;\n\tmov.b32 \t%r122, %r117;\n\tmov.b32 \t%r123, %r117;\n\tmov.b32 \t%r124, %r117;\n\tmov.b32 \t%r125, %r117;\n\tmov.b32 \t%r126, %r117;\n\tmov.b32 \t%r127, %r117;\n\tmov.b32 \t%r128, %r117;\n\tmov.b32 \t%r129, %r117;\n\tmov.b32 \t%r130, %r117;\n\tmov.b32 \t%r131, %r117;\n\tmov.b32 \t%r132, %r117;\n$L__BB0_3: // =>This Inner Loop Header: Depth=1\n\t.loc\t1 330 30 // generate_examples.py:330:30\n\tsetp.lt.s32 \t%p2, %r6, %r115;\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tmov.b32 \t%r66, 0;\n\t// begin inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tmov.u32 %r62, %r66;\n\t.loc\t1 330 23 // generate_examples.py:330:23\n\tmov.u32 %r63, %r66;\n\t.loc\t1 331 23 // generate_examples.py:331:23\n\tmov.u32 %r64, %r66;\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tmov.u32 %r65, %r66;\n\t@%p2 ld.global.v4.b32 { %r62, %r63, %r64, %r65 }, [ %rd12 + 0 ];\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tmov.u32 %r67, %r66;\n\t.loc\t1 330 23 // generate_examples.py:330:23\n\tmov.u32 %r68, %r66;\n\t.loc\t1 331 23 // generate_examples.py:331:23\n\tmov.u32 %r69, %r66;\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tmov.u32 %r70, %r66;\n\t@%p2 ld.global.v4.b32 { %r67, %r68, %r69, %r70 }, [ %rd11 + 0 ];\n\n\t// end inline asm\n\tbar.sync \t0;\n\t// begin inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tst.shared::cta.v4.b32 [ %r71 + 0 ], { %r62, %r63, %r64, %r65 };\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tst.shared::cta.v4.b32 [ %r72 + 0 ], { %r67, %r68, %r69, %r70 };\n\n\t// end inline asm\n\t.loc\t1 333 30 // generate_examples.py:333:30\n\tsetp.lt.s32 \t%p3, %r2, %r115;\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\t// begin inline asm\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tmov.u32 %r73, %r66;\n\t.loc\t1 333 23 // generate_examples.py:333:23\n\tmov.u32 %r74, %r66;\n\t.loc\t1 334 23 // generate_examples.py:334:23\n\tmov.u32 %r75, %r66;\n\t.loc\t1 335 23 // generate_examples.py:335:23\n\tmov.u32 %r76, %r66;\n\t@%p3 ld.global.v4.b32 { %r73, %r74, %r75, %r76 }, [ %rd10 + 0 ];\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tst.shared::cta.v4.b32 [ %r77 + 0 ], { %r73, %r74, %r75, %r76 };\n\n\t// end inline asm\n\t.loc\t1 329 23 // generate_examples.py:329:23\n\tbar.sync \t0;\n\tldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r78, %r79, %r80, %r81}, [%r8];\n\tldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r86, %r87, %r88, %r89}, [%r9];\n\t.loc\t1 332 23 // generate_examples.py:332:23\n\tldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%r82, %r83, %r84, %r85}, [%r10];\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\t// begin inline asm\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r117, %r118, %r119, %r120 }, { %r78, %r79, %r80, %r81 }, { %r82, %r83 }, { %r117, %r118, %r119, %r120 };\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r121, %r122, %r123, %r124 }, { %r78, %r79, %r80, %r81 }, { %r84, %r85 }, { %r121, %r122, %r123, %r124 };\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r125, %r126, %r127, %r128 }, { %r86, %r87, %r88, %r89 }, { %r82, %r83 }, { %r125, %r126, %r127, %r128 };\n\n\t// end inline asm\n\t// begin inline asm\n\t.loc\t1 335 28 // generate_examples.py:335:28\n\tmma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %r129, %r130, %r131, %r132 }, { %r86, %r87, %r88, %r89 }, { %r84, %r85 }, { %r129, %r130, %r131, %r132 };\n\n\t// end inline asm\n\t.loc\t1 336 13 // generate_examples.py:336:13\n\tadd.s64 \t%rd12, %rd12, 32;\n\tadd.s64 \t%rd11, %rd11, 32;\n\t.loc\t1 328 9 // generate_examples.py:328:9\n\tadd.s64 \t%rd10, %rd10, %rd1;\n\tadd.s32 \t%r116, %r116, -1;\n\tadd.s32 \t%r115, %r115, -16;\n\tsetp.ne.b32 \t%p4, %r116, 0;\n\t@%p4 bra \t$L__BB0_3;\n// %bb.4: // %._crit_edge.loopexit\n\t.loc\t1 338 19 // generate_examples.py:338:19\n\tcvt.rn.f16x2.f32 \t%r143, %r132, %r131;\n\tcvt.rn.f16x2.f32 \t%r142, %r130, %r129;\n\tcvt.rn.f16x2.f32 \t%r141, %r128, %r127;\n\tcvt.rn.f16x2.f32 \t%r140, %r126, %r125;\n\tcvt.rn.f16x2.f32 \t%r139, %r124, %r123;\n\tcvt.rn.f16x2.f32 \t%r138, %r122, %r121;\n\tcvt.rn.f16x2.f32 \t%r137, %r120, %r119;\n\tcvt.rn.f16x2.f32 \t%r136, %r118, %r117;\n\tbra.uni \t$L__BB0_5;\n$L__BB0_1: // %.._crit_edge_crit_edge\n\t.loc\t1 344 9 // generate_examples.py:344:9\n\tshl.b32 \t%r135, %r1, 4;\n\tand.b32 \t%r134, %r1, 4;\n\tand.b32 \t%r133, %r1, 8;\n\tmov.b32 \t%r136, 0;\n\tmov.b32 \t%r137, %r136;\n\tmov.b32 \t%r138, %r136;\n\tmov.b32 \t%r139, %r136;\n\tmov.b32 \t%r140, %r136;\n\tmov.b32 \t%r141, %r136;\n\tmov.b32 \t%r142, %r136;\n\tmov.b32 \t%r143, %r136;\n$L__BB0_5: // %._crit_edge\n\t.loc\t1 342 22 // generate_examples.py:342:22\n\tmul.lo.s32 \t%r98, %r3, %r15;\n\tmul.lo.s32 \t%r99, %r15, %r4;\n\t.loc\t1 342 18 // generate_examples.py:342:18\n\tmad.wide.s32 \t%rd7, %r98, 2, %rd4;\n\tmad.wide.s32 \t%rd8, %r99, 2, %rd4;\n\tmul.wide.s32 \t%rd9, %r7, 2;\n\tadd.s64 \t%rd5, %rd7, %rd9;\n\tadd.s64 \t%rd6, %rd8, %rd9;\n\t.loc\t1 343 19 // generate_examples.py:343:19\n\tsetp.lt.s32 \t%p7, %r3, %r11;\n\tsetp.lt.s32 \t%p8, %r4, %r11;\n\t.loc\t1 343 44 // generate_examples.py:343:44\n\tsetp.lt.s32 \t%p9, %r7, %r12;\n\t.loc\t1 343 18 // generate_examples.py:343:18\n\tand.pred \t%p5, %p7, %p9;\n\tand.pred \t%p6, %p8, %p9;\n\t.loc\t1 344 9 // generate_examples.py:344:9\n\tbar.sync \t0;\n\tand.b32 \t%r100, %r135, 176;\n\tsetp.eq.b32 \t%p10, %r134, 0;\n\tselp.b32 \t%r101, 0, 320, %p10;\n\tshl.b32 \t%r102, %r1, 2;\n\tand.b32 \t%r103, %r102, 64;\n\tor.b32 \t%r104, %r101, %r100;\n\txor.b32 \t%r105, %r104, %r103;\n\tadd.s32 \t%r106, %r114, %r105;\n\tstmatrix.sync.aligned.m8n8.x4.shared.b16 [%r106], {%r136, %r137, %r138, %r139};\n\tbar.warp.sync \t-1;\n\tshl.b32 \t%r107, %r1, 3;\n\tand.b32 \t%r108, %r107, 176;\n\tshl.b32 \t%r109, %r5, 6;\n\tsetp.eq.b32 \t%p11, %r133, 0;\n\tselp.b32 \t%r110, 0, 320, %p11;\n\tor.b32 \t%r111, %r108, %r109;\n\txor.b32 \t%r112, %r111, %r110;\n\tadd.s32 \t%r113, %r114, %r112;\n\tld.shared.v4.b32 \t{%r90, %r91, %r92, %r93}, [%r113];\n\tbar.warp.sync \t-1;\n\tstmatrix.sync.aligned.m8n8.x4.shared.b16 [%r106], {%r140, %r141, %r142, %r143};\n\tbar.warp.sync \t-1;\n\tld.shared.v4.b32 \t{%r94, %r95, %r96, %r97}, [%r113];\n\t// begin inline asm\n\t@%p5 st.global.v4.b32 [ %rd5 + 0 ], { %r90, %r91, %r92, %r93 };\n\n\t// end inline asm\n\t// begin inline asm\n\t@%p6 st.global.v4.b32 [ %rd6 + 0 ], { %r94, %r95, %r96, %r97 };\n\n\t// end inline asm\n\t.loc\t1 292 5 // generate_examples.py:292:5\n\tret;\n$L__tmp6:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.file\t2 \"/python/lib/python3.14/site-packages/triton/language/standard.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 1 // DW_CHILDREN_yes\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 2 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 0 // DW_CHILDREN_no\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 32 // DW_AT_inline\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 3 // Abbreviation Code\n.b8 46 // DW_TAG_subprogram\n.b8 1 // DW_CHILDREN_yes\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 4 // Abbreviation Code\n.b8 29 // DW_TAG_inlined_subroutine\n.b8 0 // DW_CHILDREN_no\n.b8 49 // DW_AT_abstract_origin\n.b8 19 // DW_FORM_ref4\n.b8 17 // DW_AT_low_pc\n.b8 1 // DW_FORM_addr\n.b8 18 // DW_AT_high_pc\n.b8 1 // DW_FORM_addr\n.b8 88 // DW_AT_call_file\n.b8 11 // DW_FORM_data1\n.b8 89 // DW_AT_call_line\n.b8 5 // DW_FORM_data2\n.b8 87 // DW_AT_call_column\n.b8 11 // DW_FORM_data1\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 227 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0xdc DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n.b8 2 // Abbrev [2] 0x75:0x10 DW_TAG_subprogram\n.b8 109 // DW_AT_name\n.b8 97\n.b8 116\n.b8 109\n.b8 117\n.b8 108\n.b8 95\n.b8 107\n.b8 101\n.b8 114\n.b8 110\n.b8 101\n.b8 108\n.b8 0\n.b8 1 // DW_AT_inline\n.b8 3 // Abbrev [3] 0x85:0x61 DW_TAG_subprogram\n.b64 $L__func_begin0 // DW_AT_low_pc\n.b64 $L__func_end0 // DW_AT_high_pc\n.b32 117 // DW_AT_abstract_origin\n.b8 4 // Abbrev [4] 0x9a:0x19 DW_TAG_inlined_subroutine\n.b32 117 // DW_AT_abstract_origin\n.b64 $L__tmp1 // DW_AT_low_pc\n.b64 $L__tmp2 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 56 // DW_AT_call_line\n.b8 1\n.b8 21 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xb3:0x19 DW_TAG_inlined_subroutine\n.b32 117 // DW_AT_abstract_origin\n.b64 $L__tmp2 // DW_AT_low_pc\n.b64 $L__tmp3 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 57 // DW_AT_call_line\n.b8 1\n.b8 21 // DW_AT_call_column\n.b8 4 // Abbrev [4] 0xcc:0x19 DW_TAG_inlined_subroutine\n.b32 117 // DW_AT_abstract_origin\n.b64 $L__tmp4 // DW_AT_low_pc\n.b64 $L__tmp5 // DW_AT_high_pc\n.b8 1 // DW_AT_call_file\n.b8 72 // DW_AT_call_line\n.b8 1\n.b8 27 // DW_AT_call_column\n.b8 0 // End Of Children Mark\n.b8 0 // End Of Children Mark\n\t}\n\t.section\t.debug_str\n\t{\n$L__info_string0:\n.b8 109 // string offset=0 ; matmul_kernel\n.b8 97\n.b8 116\n.b8 109\n.b8 117\n.b8 108\n.b8 95\n.b8 107\n.b8 101\n.b8 114\n.b8 110\n.b8 101\n.b8 108\n.b8 0\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","matmul_kernel.json":"{\"hash\": \"2645961e93d089721e51c9760195f3a7cabfac6742babb6be3a4dbc4331529eb\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 1, \"num_ctas\": 1, \"num_stages\": 1, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.8.0\", \"tensordesc_meta\": [], \"shared\": 1536, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"matmul_kernel\"}","matmul_kernel.sass":"Function:matmul_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 46\n /*0010*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 311\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*0020*/ S2R R7, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 99\n /*0030*/ UMOV UR4, 0x400 ;\n\t//## File \".nv_debug_ptx_txt\", line 26\n /*0040*/ ULDC.64 UR6, c[0x0][0x208] ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 95\n /*0050*/ LDC R60, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 99\n /*0060*/ S2UR UR5, SR_CgaCtaId ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \".nv_debug_ptx_txt\", line 55\n /*0070*/ VIADD R0, R9, 0xf ;\n\t//## File \".nv_debug_ptx_txt\", line 57\n /*0080*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 95\n /*0090*/ VIADD R60, R60, 0xf ;\n\t//## File \".nv_debug_ptx_txt\", line 59\n /*00a0*/ LEA.HI R3, R3, R0, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*00b0*/ IABS R10, R7 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 313\n\t//## File \".nv_debug_ptx_txt\", line 60\n /*00c0*/ SHF.R.S32.HI R4, RZ, 0x4, R3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 99\n /*00d0*/ ULEA UR4, UR5, UR4, 0x18 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*00e0*/ IABS R5, R4.reuse ;\n /*00f0*/ IABS R12, R4 ;\n /*0100*/ I2F.RP R0, R5 ;\n /*0110*/ MUFU.RCP R0, R0 ;\n /*0120*/ VIADD R2, R0, 0xffffffe ;\n /*0130*/ IMAD.MOV R0, RZ, RZ, -R12 ;\n /*0140*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0150*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0160*/ IMAD.MOV R6, RZ, RZ, -R3 ;\n /*0170*/ IMAD R11, R6, R5, RZ ;\n /*0180*/ IMAD.MOV.U32 R6, RZ, RZ, R10 ;\n /*0190*/ IMAD.HI.U32 R3, R3, R11, R2 ;\n /*01a0*/ IMAD.HI.U32 R3, R3, R6, RZ ;\n /*01b0*/ IMAD R0, R3, R0, R6 ;\n /*01c0*/ ISETP.GT.U32.AND P1, PT, R5, R0, PT ;\n /*01d0*/ @!P1 IMAD.IADD R0, R0, 0x1, -R5 ;\n /*01e0*/ @!P1 VIADD R3, R3, 0x1 ;\n /*01f0*/ ISETP.NE.AND P1, PT, R4, RZ, PT ;\n /*0200*/ ISETP.GE.U32.AND P0, PT, R0, R5, PT ;\n /*0210*/ LOP3.LUT R0, R7, R4, RZ, 0x3c, !PT ;\n /*0220*/ ISETP.GE.AND P2, PT, R0, RZ, PT ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 46\n /*0230*/ VIADD R0, R8, 0x1f ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0240*/ @P0 VIADD R3, R3, 0x1 ;\n /*0250*/ IMAD.MOV.U32 R12, RZ, RZ, R3 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*0260*/ SHF.R.S32.HI R3, RZ, 0x1f, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0270*/ @!P2 IMAD.MOV R12, RZ, RZ, -R12 ;\n\t//## File \"/python/lib/python3.14/site-packages/triton/language/standard.py\", line 43 inlined at \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 312\n\t//## File \".nv_debug_ptx_txt\", line 50\n /*0280*/ LEA.HI R3, R3, R0, RZ, 0x5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 315\n\t//## File \".nv_debug_ptx_txt\", line 63\n /*0290*/ @!P1 LOP3.LUT R12, RZ, R4, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*02a0*/ ISETP.GE.AND P2, PT, R7, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*02b0*/ LEA.HI.SX32 R3, R3, -R12, 0x1b ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*02c0*/ IMAD.MOV R5, RZ, RZ, -R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 317\n\t//## File \".nv_debug_ptx_txt\", line 67\n /*02d0*/ VIMNMX R15, R3, 0x1, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*02e0*/ IMAD R4, R4, R5, R7 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*02f0*/ IABS R11, R15.reuse ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0300*/ IABS R14, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0310*/ I2F.RP R0, R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0320*/ LOP3.LUT R4, R4, R15, RZ, 0x3c, !PT ;\n /*0330*/ ISETP.GE.AND P4, PT, R4, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0340*/ MUFU.RCP R0, R0 ;\n /*0350*/ VIADD R2, R0, 0xffffffe ;\n /*0360*/ F2I.FTZ.U32.TRUNC.NTZ R3, R2 ;\n /*0370*/ IMAD.MOV.U32 R2, RZ, RZ, RZ ;\n /*0380*/ IMAD.MOV R10, RZ, RZ, -R3 ;\n /*0390*/ IMAD R5, R10, R11, RZ ;\n /*03a0*/ IABS R10, R15 ;\n /*03b0*/ IMAD.HI.U32 R3, R3, R5, R2 ;\n /*03c0*/ IMAD.MOV R13, RZ, RZ, -R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*03d0*/ IMAD.MOV.U32 R10, RZ, RZ, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*03e0*/ IMAD.HI.U32 R0, R3, R6, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*03f0*/ IMAD.HI.U32 R5, R3, R10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0400*/ IMAD R2, R0, R13.reuse, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0410*/ IMAD R3, R5, R13, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 80\n /*0420*/ S2R R0, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0430*/ ISETP.GT.U32.AND P0, PT, R11.reuse, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0440*/ ISETP.GT.U32.AND P3, PT, R11, R3, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0450*/ @!P0 IMAD.IADD R2, R2, 0x1, -R11.reuse ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0460*/ @!P3 IMAD.IADD R3, R3, 0x1, -R11 ;\n /*0470*/ @!P3 VIADD R5, R5, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0480*/ ISETP.GT.U32.AND P1, PT, R11, R2, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0490*/ ISETP.GE.U32.AND P0, PT, R3, R11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*04a0*/ LOP3.LUT R3, RZ, R15, RZ, 0x33, !PT ;\n /*04b0*/ @!P1 IMAD.IADD R2, R2, 0x1, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*04c0*/ @P0 VIADD R5, R5, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*04d0*/ ISETP.NE.AND P0, PT, R15, RZ, PT ;\n /*04e0*/ @!P2 IMAD.MOV R2, RZ, RZ, -R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*04f0*/ @!P4 IMAD.MOV R5, RZ, RZ, -R5 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0500*/ SEL R2, R3.reuse, R2, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 319\n\t//## File \".nv_debug_ptx_txt\", line 76\n /*0510*/ SEL R7, R3, R5, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 98\n /*0520*/ ISETP.GT.AND P0, PT, R60, 0xf, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 318\n\t//## File \".nv_debug_ptx_txt\", line 71\n /*0530*/ IMAD.IADD R6, R12, 0x1, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*0540*/ SHF.R.U32.HI R3, RZ, 0x1, R0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 88\n /*0550*/ LOP3.LUT R2, R0, 0x1, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*0560*/ IMAD.SHL.U32 R6, R6, 0x20, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*0570*/ SGXT.U32 R3, R3, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 89\n /*0580*/ IMAD.SHL.U32 R4, R2, 0x8, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 83\n /*0590*/ LOP3.LUT R5, R6, R3, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 91\n /*05a0*/ IMAD R6, R7, 0x10, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 84\n /*05b0*/ LOP3.LUT R7, R5, 0x10, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 100\n /*05c0*/ @!P0 BRA `(.L_x_0) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*05d0*/ IABS R17, R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*05e0*/ LDC R61, c[0x0][0x238] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*05f0*/ IABS R16, R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*0600*/ ULDC UR5, c[0x0][0x234] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0610*/ I2F.RP R14, R17 ;\n /*0620*/ ISETP.NE.AND P3, PT, R9, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0630*/ ULDC.64 UR8, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 134\n /*0640*/ SHF.R.S32.HI R23, RZ, 0x3, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0650*/ ULDC.64 UR10, c[0x0][0x210] ;\n\t//## File \".nv_debug_ptx_txt\", line 131\n /*0660*/ IMAD.SHL.U32 R56, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 133\n /*0670*/ LOP3.LUT R58, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0680*/ I2F.RP R15, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 134\n /*0690*/ SGXT R23, R23, 0x1 ;\n\t//## File \".nv_debug_ptx_txt\", line 142\n /*06a0*/ LOP3.LUT R59, R0, 0x4, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 135\n /*06b0*/ LOP3.LUT R25, R23, 0x90, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*06c0*/ CS2R R22, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*06d0*/ MUFU.RCP R14, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 136\n /*06e0*/ LOP3.LUT R48, R25, 0x170, R56, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*06f0*/ MUFU.RCP R15, R15 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0700*/ VIADD R10, R14, 0xffffffe ;\n /*0710*/ IABS R14, R6 ;\n /*0720*/ F2I.FTZ.U32.TRUNC.NTZ R11, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0730*/ VIADD R12, R15, 0xffffffe ;\n /*0740*/ F2I.FTZ.U32.TRUNC.NTZ R13, R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0750*/ IMAD.MOV.U32 R10, RZ, RZ, RZ ;\n /*0760*/ IMAD.MOV R18, RZ, RZ, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0770*/ IMAD.MOV.U32 R12, RZ, RZ, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0780*/ IMAD R19, R18, R17, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0790*/ IMAD.MOV R21, RZ, RZ, -R13 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*07a0*/ IMAD.HI.U32 R11, R11, R19, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*07b0*/ IABS R10, R7 ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*07c0*/ IABS R19, R5 ;\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*07d0*/ IMAD R15, R21, R16, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*07e0*/ CS2R R20, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*07f0*/ IMAD.HI.U32 R11, R11, R14, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0800*/ IMAD.HI.U32 R12, R13, R15, R12 ;\n /*0810*/ IMAD.MOV.U32 R13, RZ, RZ, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0820*/ IMAD.MOV R11, RZ, RZ, -R11 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0830*/ IMAD.HI.U32 R10, R12, R13, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0840*/ IMAD.HI.U32 R12, R12, R19, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0850*/ IMAD.MOV R15, RZ, RZ, -R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0860*/ IMAD.MOV R18, RZ, RZ, -R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0870*/ IMAD R12, R16.reuse, R15, R13 ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0880*/ IMAD R13, R16.reuse, R18, R19 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0890*/ CS2R R18, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*08a0*/ IMAD R10, R17, R11, R14 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*08b0*/ ISETP.GT.U32.AND P1, PT, R16.reuse, R12, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*08c0*/ ISETP.GT.U32.AND P2, PT, R16, R13, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*08d0*/ ISETP.GT.U32.AND P0, PT, R17, R10, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 107\n /*08e0*/ SHF.R.S32.HI R11, RZ, 0x1f, R60 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 143\n /*08f0*/ SHF.R.S32.HI R14, RZ, 0x2, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 109\n /*0900*/ LEA.HI R60, R11, R60, RZ, 0x4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0910*/ LOP3.LUT R11, RZ, R8, RZ, 0x33, !PT ;\n /*0920*/ @!P1 IMAD.IADD R12, R12, 0x1, -R16.reuse ;\n /*0930*/ ISETP.GE.AND P1, PT, R7, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0940*/ @!P2 IMAD.IADD R13, R13, 0x1, -R16 ;\n /*0950*/ ISETP.GE.AND P2, PT, R5, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0960*/ @!P0 IMAD.IADD R10, R10, 0x1, -R17 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0970*/ ISETP.GT.U32.AND P5, PT, R16, R12, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0980*/ ISETP.GT.U32.AND P6, PT, R16, R13, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0990*/ ISETP.GT.U32.AND P4, PT, R17, R10, PT ;\n /*09a0*/ ISETP.GE.AND P0, PT, R6, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 110\n /*09b0*/ SHF.R.S32.HI R60, RZ, 0x4, R60 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*09c0*/ @!P5 IMAD.IADD R12, R12, 0x1, -R16.reuse ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*09d0*/ @!P6 IMAD.IADD R13, R13, 0x1, -R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*09e0*/ @!P4 IMAD.IADD R10, R10, 0x1, -R17 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*09f0*/ ISETP.NE.AND P4, PT, R8, RZ, PT ;\n /*0a00*/ @!P1 IMAD.MOV R12, RZ, RZ, -R12 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0a10*/ CS2R R16, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0a20*/ @!P2 IMAD.MOV R13, RZ, RZ, -R13 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 322\n\t//## File \".nv_debug_ptx_txt\", line 114\n /*0a30*/ @!P0 IMAD.MOV R10, RZ, RZ, -R10 ;\n /*0a40*/ @!P3 LOP3.LUT R10, RZ, R9, RZ, 0x33, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 140\n /*0a50*/ IMAD.SHL.U32 R8, R0, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 321\n\t//## File \".nv_debug_ptx_txt\", line 120\n /*0a60*/ SEL R9, R11, R12, !P4 ;\n\t//## File \".nv_debug_ptx_txt\", line 126\n /*0a70*/ SEL R11, R11, R13, !P4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 116\n /*0a80*/ IMAD R10, R3, R61, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 143\n /*0a90*/ SGXT R12, R14, 0x1 ;\n\t//## File \".nv_debug_ptx_txt\", line 122\n /*0aa0*/ IMAD R9, R9, UR5, R4.reuse ;\n\t//## File \".nv_debug_ptx_txt\", line 141\n /*0ab0*/ LOP3.LUT R8, R8, 0x160, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 128\n /*0ac0*/ IMAD R11, R11, UR5, R4 ;\n\t//## File \".nv_debug_ptx_txt\", line 144\n /*0ad0*/ LOP3.LUT R13, R12, 0x90, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0ae0*/ CS2R R14, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0af0*/ LEA R50, P1, R9, UR10, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 337\n\t//## File \".nv_debug_ptx_txt\", line 112\n /*0b00*/ IMAD.SHL.U32 R61, R61, 0x10, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 150\n /*0b10*/ LOP3.LUT R49, R13, 0x10, R0, 0x78, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 39\n /*0b20*/ ULDC UR5, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 130\n /*0b30*/ LEA R54, P2, R11, UR10, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0b40*/ LEA R38, P0, R10, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 146\n /*0b50*/ LOP3.LUT R51, R13, R8, RZ, 0xfc, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0b60*/ CS2R R12, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 152\n /*0b70*/ IADD3 R49, R49, UR4, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 124\n /*0b80*/ LEA.HI.X.SX32 R55, R9, UR11, 0x1, P1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0b90*/ CS2R R8, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 130\n /*0ba0*/ LEA.HI.X.SX32 R57, R11, UR11, 0x1, P2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 325\n\t//## File \".nv_debug_ptx_txt\", line 118\n /*0bb0*/ LEA.HI.X.SX32 R39, R10, UR9, 0x1, P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 171\n /*0bc0*/ CS2R R10, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 324\n\t//## File \".nv_debug_ptx_txt\", line 147\n /*0bd0*/ LOP3.LUT R51, R51, 0x10, R0, 0x78, !PT ;\n.L_x_1:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 330\n\t//## File \".nv_debug_ptx_txt\", line 174\n /*0be0*/ ISETP.GE.AND P1, PT, R4, UR5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 186\n /*0bf0*/ CS2R R24, SRZ ;\n /*0c00*/ CS2R R26, SRZ ;\n\t//## File \".nv_debug_ptx_txt\", line 198\n /*0c10*/ CS2R R28, SRZ ;\n /*0c20*/ CS2R R30, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 333\n\t//## File \".nv_debug_ptx_txt\", line 213\n /*0c30*/ ISETP.GE.AND P0, PT, R3, UR5, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 186\n /*0c40*/ @!P1 IMAD.MOV.U32 R36, RZ, RZ, R54 ;\n /*0c50*/ @!P1 IMAD.MOV.U32 R37, RZ, RZ, R57 ;\n\t//## File \".nv_debug_ptx_txt\", line 198\n /*0c60*/ @!P1 IMAD.MOV.U32 R52, RZ, RZ, R50 ;\n /*0c70*/ @!P1 IMAD.MOV.U32 R53, RZ, RZ, R55 ;\n\t//## File \".nv_debug_ptx_txt\", line 186\n /*0c80*/ @!P1 LDG.E.128 R24, desc[UR6][R36.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 224\n /*0c90*/ CS2R R32, SRZ ;\n /*0ca0*/ CS2R R34, SRZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 198\n /*0cb0*/ @!P1 LDG.E.128 R28, desc[UR6][R52.64] ;\n\t//## File \".nv_debug_ptx_txt\", line 201\n /*0cc0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n /*0cd0*/ IMAD.MOV.U32 R52, RZ, RZ, R38 ;\n /*0ce0*/ IMAD.MOV.U32 R53, RZ, RZ, R39 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 224\n /*0cf0*/ @!P0 LDG.E.128 R32, desc[UR6][R52.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 264\n /*0d00*/ VIADD R60, R60, 0xffffffff ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 260\n /*0d10*/ IADD3 R54, P1, R54, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 265\n /*0d20*/ UIADD3 UR5, UR5, -0x10, URZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 261\n /*0d30*/ IADD3 R50, P2, R50, 0x20, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 266\n /*0d40*/ ISETP.NE.U32.AND P0, PT, R60, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 336\n\t//## File \".nv_debug_ptx_txt\", line 260\n /*0d50*/ IMAD.X R57, RZ, RZ, R57, P1 ;\n\t//## File \".nv_debug_ptx_txt\", line 261\n /*0d60*/ IMAD.X R55, RZ, RZ, R55, P2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 329\n\t//## File \".nv_debug_ptx_txt\", line 204\n /*0d70*/ STS.128 [R48+UR4], R24 ;\n\t//## File \".nv_debug_ptx_txt\", line 209\n /*0d80*/ STS.128 [R48+UR4+0x200], R28 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 229\n /*0d90*/ STS.128 [R48+UR4+0x400], R32 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 329\n\t//## File \".nv_debug_ptx_txt\", line 233\n /*0da0*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 234\n /*0db0*/ LDSM.16.M88.4 R36, [R51+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 332\n\t//## File \".nv_debug_ptx_txt\", line 237\n /*0dc0*/ LDSM.16.MT88.4 R40, [R49+0x400] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 329\n\t//## File \".nv_debug_ptx_txt\", line 235\n /*0dd0*/ LDSM.16.M88.4 R44, [R51+UR4+0x200] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 241\n /*0de0*/ HMMA.16816.F32 R8, R36.reuse, R40, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 246\n /*0df0*/ HMMA.16816.F32 R12, R36, R42, R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 251\n /*0e00*/ HMMA.16816.F32 R16, R44, R40, R16 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 263\n /*0e10*/ IMAD.WIDE R38, R61, 0x2, R52 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 335\n\t//## File \".nv_debug_ptx_txt\", line 256\n /*0e20*/ HMMA.16816.F32 R20, R44, R42, R20 ;\n /*0e30*/ NOP ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 328\n\t//## File \".nv_debug_ptx_txt\", line 267\n /*0e40*/ @P0 BRA `(.L_x_1) ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 338\n\t//## File \".nv_debug_ptx_txt\", line 270\n /*0e50*/ F2FP.F16.F32.PACK_AB R23, R23, R22 ;\n\t//## File \".nv_debug_ptx_txt\", line 274\n /*0e60*/ F2FP.F16.F32.PACK_AB R15, R15, R14 ;\n\t//## File \".nv_debug_ptx_txt\", line 271\n /*0e70*/ F2FP.F16.F32.PACK_AB R22, R21, R20 ;\n\t//## File \".nv_debug_ptx_txt\", line 275\n /*0e80*/ F2FP.F16.F32.PACK_AB R14, R13, R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 272\n /*0e90*/ F2FP.F16.F32.PACK_AB R21, R19, R18 ;\n\t//## File \".nv_debug_ptx_txt\", line 273\n /*0ea0*/ F2FP.F16.F32.PACK_AB R20, R17, R16 ;\n\t//## File \".nv_debug_ptx_txt\", line 276\n /*0eb0*/ F2FP.F16.F32.PACK_AB R13, R11, R10 ;\n\t//## File \".nv_debug_ptx_txt\", line 277\n /*0ec0*/ F2FP.F16.F32.PACK_AB R12, R9, R8 ;\n\t//## File \".nv_debug_ptx_txt\", line 278\n /*0ed0*/ BRA `(.L_x_2) ;\n.L_x_0:\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 281\n /*0ee0*/ IMAD.SHL.U32 R56, R0.reuse, 0x10, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 282\n /*0ef0*/ LOP3.LUT R59, R0.reuse, 0x4, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 291\n /*0f00*/ CS2R R12, SRZ ;\n\t//## File \".nv_debug_ptx_txt\", line 283\n /*0f10*/ LOP3.LUT R58, R0, 0x8, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 291\n /*0f20*/ CS2R R14, SRZ ;\n /*0f30*/ CS2R R20, SRZ ;\n /*0f40*/ CS2R R22, SRZ ;\n.L_x_2:\n\t//## File \".nv_debug_ptx_txt\", line 311\n /*0f50*/ BAR.SYNC.DEFER_BLOCKING 0x0 ;\n\t//## File \".nv_debug_ptx_txt\", line 313\n /*0f60*/ ISETP.NE.U32.AND P0, PT, R59, RZ, PT ;\n\t//## File \".nv_debug_ptx_txt\", line 315\n /*0f70*/ IMAD.SHL.U32 R4, R0.reuse, 0x4, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 322\n /*0f80*/ IMAD.SHL.U32 R0, R0, 0x8, RZ ;\n\t//## File \".nv_debug_ptx_txt\", line 314\n /*0f90*/ SEL R3, RZ, 0x140, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 294\n /*0fa0*/ ULDC UR5, c[0x0][0x23c] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 325\n /*0fb0*/ ISETP.NE.U32.AND P0, PT, R58, RZ, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 306\n /*0fc0*/ ULDC.64 UR10, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 317\n /*0fd0*/ LOP3.LUT R3, R3, 0xb0, R56, 0xf8, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 297\n /*0fe0*/ ULDC.64 UR8, c[0x0][0x220] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 323\n /*0ff0*/ LOP3.LUT R9, R0, 0xb0, RZ, 0xc0, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 295\n /*1000*/ IMAD R8, R7, UR5, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 318\n /*1010*/ LOP3.LUT R3, R3, 0x40, R4, 0x78, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 326\n /*1020*/ SEL R0, RZ, 0x140, !P0 ;\n\t//## File \".nv_debug_ptx_txt\", line 327\n /*1030*/ IMAD R11, R2, 0x40, R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 343\n\t//## File \".nv_debug_ptx_txt\", line 306\n /*1040*/ ISETP.GE.AND P0, PT, R6, UR11, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 319\n /*1050*/ VIADD R9, R3, UR4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 298\n /*1060*/ LEA R4, P3, R8, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 328\n /*1070*/ LOP3.LUT R11, R11, R0, RZ, 0x3c, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 294\n /*1080*/ IMAD R0, R5.reuse, UR5, RZ ;\n /*1090*/ ISETP.LT.AND P1, PT, R5, UR10, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 320\n /*10a0*/ STSM.16.M88.4 [R9], R12 ;\n\t//## File \".nv_debug_ptx_txt\", line 321\n /*10b0*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 330\n /*10c0*/ LDS.128 R16, [R11+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 297\n /*10d0*/ LEA R2, P2, R0.reuse, UR8, 0x1 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 331\n /*10e0*/ NOP ;\n /*10f0*/ ISETP.LT.AND P0, PT, R7, UR10, !P0 ;\n\t//## File \".nv_debug_ptx_txt\", line 332\n /*1100*/ STSM.16.M88.4 [R9], R20 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 297\n /*1110*/ LEA.HI.X.SX32 R3, R0, UR9, 0x1, P2 ;\n\t//## File \".nv_debug_ptx_txt\", line 298\n /*1120*/ LEA.HI.X.SX32 R5, R8, UR9, 0x1, P3 ;\n\t//## File \".nv_debug_ptx_txt\", line 300\n /*1130*/ IMAD.WIDE R2, R6, 0x2, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 333\n /*1140*/ NOP ;\n\t//## File \".nv_debug_ptx_txt\", line 336\n /*1150*/ @P1 STG.E.128 desc[UR6][R2.64], R16 ;\n\t//## File \".nv_debug_ptx_txt\", line 340\n /*1160*/ @!P0 EXIT ;\n\t//## File \".nv_debug_ptx_txt\", line 334\n /*1170*/ LDS.128 R8, [R11+UR4] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 342\n\t//## File \".nv_debug_ptx_txt\", line 301\n /*1180*/ IMAD.WIDE R4, R6, 0x2, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 344\n\t//## File \".nv_debug_ptx_txt\", line 340\n /*1190*/ STG.E.128 desc[UR6][R4.64], R8 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 292\n\t//## File \".nv_debug_ptx_txt\", line 344\n /*11a0*/ EXIT ;\n.L_x_3:\n /*11b0*/ BRA `(.L_x_3);\n /*11c0*/ NOP;\n /*11d0*/ NOP;\n /*11e0*/ NOP;\n /*11f0*/ NOP;\n /*1200*/ NOP;\n /*1210*/ NOP;\n /*1220*/ NOP;\n /*1230*/ NOP;\n /*1240*/ NOP;\n /*1250*/ NOP;\n /*1260*/ NOP;\n /*1270*/ NOP;\n.L_x_4:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":266,"end_line":344,"code":" @triton.autotune(\n configs=[\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 16,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n triton.Config(\n {\n \"BLOCK_SIZE_M\": 32,\n \"BLOCK_SIZE_N\": 16,\n \"BLOCK_SIZE_K\": 16,\n \"GROUP_SIZE_M\": 1,\n },\n num_stages=1,\n num_warps=1,\n ),\n ],\n key=[\"M\", \"N\", \"K\"],\n )\n @triton.jit\n def matmul_kernel(\n a,\n b,\n c,\n M,\n N,\n K,\n stride_am,\n stride_ak,\n stride_bk,\n stride_bn,\n stride_cm,\n stride_cn,\n BLOCK_SIZE_M: tl.constexpr,\n BLOCK_SIZE_N: tl.constexpr,\n BLOCK_SIZE_K: tl.constexpr,\n GROUP_SIZE_M: tl.constexpr,\n ACTIVATION: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)\n num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)\n num_pid_in_group = GROUP_SIZE_M * num_pid_n\n group_id = pid // num_pid_in_group\n first_pid_m = group_id * GROUP_SIZE_M\n group_size = min(num_pid_m - first_pid_m, GROUP_SIZE_M)\n pid_m = first_pid_m + (pid % group_size)\n pid_n = (pid % num_pid_in_group) // group_size\n\n offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M\n offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N\n offs_k = tl.arange(0, BLOCK_SIZE_K)\n a_ptrs = a + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)\n b_ptrs = b + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)\n\n accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)\n for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):\n a_block = tl.load(\n a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0\n )\n b_block = tl.load(\n b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0\n )\n accumulator += tl.dot(a_block, b_block)\n a_ptrs += BLOCK_SIZE_K * stride_ak\n b_ptrs += BLOCK_SIZE_K * stride_bk\n c_block = accumulator.to(tl.float16)\n\n offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)\n offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)\n c_ptrs = c + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]\n c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)\n tl.store(c_ptrs, c_block, mask=c_mask)\n"}}} +{"event_type":"autotune","pid":4104992,"kernel_name":"matmul_kernel","cache_key":"528f8911847b1f39f462157322ec9a674a5565b9b74d621d5ff09625eb710f38","best_config":"BLOCK_SIZE_M: 32, BLOCK_SIZE_N: 16, BLOCK_SIZE_K: 16, GROUP_SIZE_M: 1, num_warps: 1, num_ctas: 1, num_stages: 1, maxnreg: None","configs_timings":{"BLOCK_SIZE_M: 16, BLOCK_SIZE_N: 16, BLOCK_SIZE_K: 16, GROUP_SIZE_M: 1, num_warps: 1, num_ctas: 1, num_stages: 1, maxnreg: None":[0.25833600759506226,0.24760960042476654,0.26989439129829407],"BLOCK_SIZE_M: 32, BLOCK_SIZE_N: 16, BLOCK_SIZE_K: 16, GROUP_SIZE_M: 1, num_warps: 1, num_ctas: 1, num_stages: 1, maxnreg: None":[0.24879999458789825,0.23439359664916992,0.431852787733078]},"duration":1.7164421081542969,"cache_hit":false,"autotune_key":"(16, 16, 16, 'torch.float16', 'torch.float16', 'torch.float16')","stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":257,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"knobs.autotuning.listener("},{"line":2114,"name":"_autotune_listener","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"autotune\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"autotune","pid":4104992,"kernel_name":"matmul_kernel","cache_key":"528f8911847b1f39f462157322ec9a674a5565b9b74d621d5ff09625eb710f38","best_config":"BLOCK_SIZE_M: 16, BLOCK_SIZE_N: 16, BLOCK_SIZE_K: 16, GROUP_SIZE_M: 1, num_warps: 1, num_ctas: 1, num_stages: 1, maxnreg: None","configs_timings":{"BLOCK_SIZE_M: 16, BLOCK_SIZE_N: 16, BLOCK_SIZE_K: 16, GROUP_SIZE_M: 1, num_warps: 1, num_ctas: 1, num_stages: 1, maxnreg: None":[0.2383359968662262,0.22753280103206636,0.24840319454669954],"BLOCK_SIZE_M: 32, BLOCK_SIZE_N: 16, BLOCK_SIZE_K: 16, GROUP_SIZE_M: 1, num_warps: 1, num_ctas: 1, num_stages: 1, maxnreg: None":[0.2393600046634674,0.22932479679584503,0.24669439494609832]},"duration":0.24884319305419922,"cache_hit":false,"autotune_key":"(32, 32, 16, 'torch.float16', 'torch.float16', 'torch.float16')","stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":421,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul(a, b).sum()"},{"line":382,"name":"matmul","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"matmul_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":257,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/autotuner.py","loc":"knobs.autotuning.listener("},{"line":2114,"name":"_autotune_listener","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton(\"autotune\", metadata_fn=lambda: convert(trace_data))"},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:04.%fZ"} +{"event_type":"compilation","pid":4104992,"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":429,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op(x, y, z, scale_factor=1.0, activation=\"none\").sum()"},{"line":404,"name":"fused_op","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":757,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":902,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:05.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"0a7707c5f05286447121bb3674a0dbe36368322c817f490488942e13fc939971","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"fused_op_kernel","num_ctas":1,"num_stages":3,"num_warps":4,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":0,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[],"(6,)":[],"(7,)":[]},"src_cache_key":"f2e5ae6c41d3bfc26af60eada8fff38a0b318cffe77df20b32d117346110ff07","src_constants":{"(6,)":"none","(7,)":8},"times":{"ir_initialization":1641,"lowering_stages":[],"store_results":0}},"file_path":{"fused_op_kernel.source":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.source","fused_op_kernel.ttir":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.ttir","fused_op_kernel.ttgir":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.ttgir","fused_op_kernel.llir":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.llir","fused_op_kernel.ptx":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.ptx","fused_op_kernel.cubin":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.cubin","fused_op_kernel.json":"/home/user/.triton/cache/BJ3QPRPQKKDEI4JBXM3HJIG34NRWQMRMQF7USBEISQXBH7ETTFYQ/fused_op_kernel.json"},"file_content":{"fused_op_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:5)\n#loc15 = loc(\"a_ptr\"(#loc))\n#loc16 = loc(\"b_ptr\"(#loc))\n#loc17 = loc(\"c_ptr\"(#loc))\n#loc18 = loc(\"output_ptr\"(#loc))\n#loc19 = loc(\"n_elements\"(#loc))\n#loc20 = loc(\"scale_factor\"(#loc))\nmodule {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc21)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc22)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc23)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32> loc(#loc22)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32> loc(#loc22)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32> loc(#loc24)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32> loc(#loc24)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc25)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc25)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr> loc(#loc26)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc27)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc27)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr> loc(#loc28)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc29)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc29)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr> loc(#loc30)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32> loc(#loc31)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32> loc(#loc31)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32> loc(#loc31)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32> loc(#loc31)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc13)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc13)\n tt.store %1, %result_12, %mask_3 : tensor<8x!tt.ptr> loc(#loc14)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:15)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:19)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:38)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:16)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:13)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:21)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:13)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:21)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:13)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:18)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:18)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:9)\n#loc21 = loc(\"pid\"(#loc2))\n#loc22 = loc(\"offsets\"(#loc3))\n#loc23 = loc(\"offsets\"(#loc4))\n#loc24 = loc(\"mask\"(#loc5))\n#loc25 = loc(\"a\"(#loc6))\n#loc26 = loc(\"a\"(#loc7))\n#loc27 = loc(\"b\"(#loc8))\n#loc28 = loc(\"b\"(#loc9))\n#loc29 = loc(\"c\"(#loc10))\n#loc30 = loc(\"c\"(#loc11))\n#loc31 = loc(\"result\"(#loc12))\n","fused_op_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:5)\n#loc15 = loc(\"a_ptr\"(#loc))\n#loc16 = loc(\"b_ptr\"(#loc))\n#loc17 = loc(\"c_ptr\"(#loc))\n#loc18 = loc(\"output_ptr\"(#loc))\n#loc19 = loc(\"n_elements\"(#loc))\n#loc20 = loc(\"scale_factor\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 4 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc21)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc22)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #blocked> loc(#loc23)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32, #blocked> loc(#loc22)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32, #blocked> loc(#loc22)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32, #blocked> loc(#loc24)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32, #blocked> loc(#loc24)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc25)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc25)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc26)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc27)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc27)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc28)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc29)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc29)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc30)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32, #blocked> loc(#loc31)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32, #blocked> loc(#loc31)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32, #blocked> loc(#loc31)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32, #blocked> loc(#loc31)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc13)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc13)\n tt.store %1, %result_12, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc14)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:15)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:19)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:38)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:16)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:13)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:21)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:13)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:21)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:13)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:18)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:18)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:9)\n#loc21 = loc(\"pid\"(#loc2))\n#loc22 = loc(\"offsets\"(#loc3))\n#loc23 = loc(\"offsets\"(#loc4))\n#loc24 = loc(\"mask\"(#loc5))\n#loc25 = loc(\"a\"(#loc6))\n#loc26 = loc(\"a\"(#loc7))\n#loc27 = loc(\"b\"(#loc8))\n#loc28 = loc(\"b\"(#loc9))\n#loc29 = loc(\"c\"(#loc10))\n#loc30 = loc(\"c\"(#loc11))\n#loc31 = loc(\"result\"(#loc12))\n","fused_op_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @fused_op_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, i32 %4, float %5, ptr addrspace(1) readnone captures(none) %6, ptr addrspace(1) readnone captures(none) %7) local_unnamed_addr #0 !dbg !4 {\n %9 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %10 = shl i32 %9, 3, !dbg !13\n %11 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !14\n %12 = and i32 %11, 7, !dbg !14\n %13 = or disjoint i32 %10, %12, !dbg !13\n %14 = icmp slt i32 %13, %4, !dbg !15\n %15 = sext i32 %13 to i64, !dbg !16\n %16 = getelementptr [4 x i8], ptr addrspace(1) %0, i64 %15, !dbg !16\n %17 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %16, i1 %14) #2, !dbg !17\n %18 = bitcast i32 %17 to float, !dbg !17\n %19 = getelementptr [4 x i8], ptr addrspace(1) %1, i64 %15, !dbg !18\n %20 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %19, i1 %14) #2, !dbg !19\n %21 = bitcast i32 %20 to float, !dbg !19\n %22 = getelementptr [4 x i8], ptr addrspace(1) %2, i64 %15, !dbg !20\n %23 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %22, i1 %14) #2, !dbg !21\n %24 = bitcast i32 %23 to float, !dbg !21\n %25 = fmul float %18, %21, !dbg !22\n %26 = fmul float %5, %25, !dbg !22\n %27 = fadd float %26, %24, !dbg !22\n %28 = getelementptr [4 x i8], ptr addrspace(1) %3, i64 %15, !dbg !23\n %29 = and i32 %11, 120, !dbg !24\n %30 = icmp eq i32 %29, 0, !dbg !24\n %31 = bitcast float %27 to i32, !dbg !24\n %32 = and i1 %30, %14, !dbg !24\n tail call void asm sideeffect \"@$2 st.global.b32 [ $1 + 0 ], { $0 };\", \"r,l,b\"(i32 %31, ptr addrspace(1) %28, i1 %32) #2, !dbg !24\n ret void, !dbg !25\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"128\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"fused_op_kernel\", linkageName: \"fused_op_kernel\", scope: !1, file: !1, line: 347, type: !5, scopeLine: 347, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !7, !9, !8, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"float\", size: 32, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 357, column: 15, scope: !4)\n!13 = !DILocation(line: 358, column: 19, scope: !4)\n!14 = !DILocation(line: 358, column: 38, scope: !4)\n!15 = !DILocation(line: 359, column: 16, scope: !4)\n!16 = !DILocation(line: 361, column: 21, scope: !4)\n!17 = !DILocation(line: 361, column: 13, scope: !4)\n!18 = !DILocation(line: 362, column: 21, scope: !4)\n!19 = !DILocation(line: 362, column: 13, scope: !4)\n!20 = !DILocation(line: 363, column: 21, scope: !4)\n!21 = !DILocation(line: 363, column: 13, scope: !4)\n!22 = !DILocation(line: 365, column: 18, scope: !4)\n!23 = !DILocation(line: 369, column: 18, scope: !4)\n!24 = !DILocation(line: 369, column: 9, scope: !4)\n!25 = !DILocation(line: 347, column: 5, scope: !4)\n","fused_op_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.8\n.target sm_90a\n.address_size 64\n\n\t// .globl\tfused_op_kernel // -- Begin function fused_op_kernel\n // @fused_op_kernel\n.visible .entry fused_op_kernel(\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_2,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_3,\n\t.param .u32 fused_op_kernel_param_4,\n\t.param .f32 fused_op_kernel_param_5,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_6,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_7\n)\n.reqntid 128\n{\n\t.reg .pred \t%p<4>;\n\t.reg .b32 \t%r<14>;\n\t.reg .b64 \t%rd<10>;\n\t.loc\t1 347 0 // generate_examples.py:347:0\n$L__func_begin0:\n\t.loc\t1 347 0 // generate_examples.py:347:0\n\n// %bb.0:\n\tld.param.b64 \t%rd5, [fused_op_kernel_param_0];\n\tld.param.b64 \t%rd6, [fused_op_kernel_param_1];\n$L__tmp0:\n\t.loc\t1 357 15 // generate_examples.py:357:15\n\tmov.u32 \t%r5, %ctaid.x;\n\t.loc\t1 358 19 // generate_examples.py:358:19\n\tshl.b32 \t%r6, %r5, 3;\n\tld.param.b64 \t%rd7, [fused_op_kernel_param_2];\n\tld.param.b64 \t%rd8, [fused_op_kernel_param_3];\n\t.loc\t1 358 38 // generate_examples.py:358:38\n\tmov.u32 \t%r7, %tid.x;\n\tand.b32 \t%r8, %r7, 7;\n\tld.param.b32 \t%r9, [fused_op_kernel_param_4];\n\t.loc\t1 358 19 // generate_examples.py:358:19\n\tor.b32 \t%r10, %r6, %r8;\n\tld.param.b32 \t%r11, [fused_op_kernel_param_5];\n\t.loc\t1 359 16 // generate_examples.py:359:16\n\tsetp.lt.s32 \t%p1, %r10, %r9;\n\t.loc\t1 361 21 // generate_examples.py:361:21\n\tmul.wide.s32 \t%rd9, %r10, 4;\n\tadd.s64 \t%rd1, %rd5, %rd9;\n\t.loc\t1 361 13 // generate_examples.py:361:13\n\t// begin inline asm\n\t.loc\t1 361 13 // generate_examples.py:361:13\n\tmov.u32 %r1, 0x0;\n\t@%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];\n\n\t// end inline asm\n\t.loc\t1 362 21 // generate_examples.py:362:21\n\tadd.s64 \t%rd2, %rd6, %rd9;\n\t.loc\t1 362 13 // generate_examples.py:362:13\n\t// begin inline asm\n\t.loc\t1 362 13 // generate_examples.py:362:13\n\tmov.u32 %r2, 0x0;\n\t@%p1 ld.global.b32 { %r2 }, [ %rd2 + 0 ];\n\n\t// end inline asm\n\t.loc\t1 363 21 // generate_examples.py:363:21\n\tadd.s64 \t%rd3, %rd7, %rd9;\n\t.loc\t1 363 13 // generate_examples.py:363:13\n\t// begin inline asm\n\t.loc\t1 363 13 // generate_examples.py:363:13\n\tmov.u32 %r3, 0x0;\n\t@%p1 ld.global.b32 { %r3 }, [ %rd3 + 0 ];\n\n\t// end inline asm\n\t.loc\t1 365 18 // generate_examples.py:365:18\n\tmul.f32 \t%r12, %r1, %r2;\n\tfma.rn.f32 \t%r4, %r11, %r12, %r3;\n\t.loc\t1 369 18 // generate_examples.py:369:18\n\tadd.s64 \t%rd4, %rd8, %rd9;\n\t.loc\t1 369 9 // generate_examples.py:369:9\n\tand.b32 \t%r13, %r7, 120;\n\tsetp.eq.b32 \t%p3, %r13, 0;\n\tand.pred \t%p2, %p3, %p1;\n\t// begin inline asm\n\t@%p2 st.global.b32 [ %rd4 + 0 ], { %r4 };\n\n\t// end inline asm\n\t.loc\t1 347 5 // generate_examples.py:347:5\n\tret;\n$L__tmp1:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 0 // DW_CHILDREN_no\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 113 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0x6a DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","fused_op_kernel.json":"{\"hash\": \"0a7707c5f05286447121bb3674a0dbe36368322c817f490488942e13fc939971\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 4, \"num_ctas\": 1, \"num_stages\": 3, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.8.0\", \"tensordesc_meta\": [], \"shared\": 0, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"fused_op_kernel\"}","fused_op_kernel.sass":"Function:fused_op_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 347\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 41\n /*0010*/ S2R R12, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*0020*/ LDC.64 R2, c[0x0][0x210] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*0030*/ ULDC UR6, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0040*/ MOV R13, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 357\n\t//## File \".nv_debug_ptx_txt\", line 35\n /*0050*/ S2R R11, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 60\n /*0060*/ LDC.64 R4, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0070*/ LDC.64 R6, c[0x0][0x220] ;\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0080*/ HFMA2.MMA R10, -RZ, RZ, 0, 0 ;\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0090*/ ULDC.64 UR4, c[0x0][0x208] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*00a0*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*00b0*/ LOP3.LUT R0, R12, 0x7, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*00c0*/ LEA R11, R11, R0, 0x3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 56\n /*00d0*/ HFMA2.MMA R0, -RZ, RZ, 0, 0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*00e0*/ ISETP.GE.AND P0, PT, R11.reuse, UR6, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*00f0*/ IMAD.WIDE R2, R11, 0x4, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 60\n /*0100*/ IMAD.WIDE R4, R11, 0x4, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0110*/ IMAD.WIDE R6, R11, 0x4, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 56\n /*0120*/ @!P0 LDG.E R0, desc[UR4][R2.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0130*/ @!P0 LDG.E R13, desc[UR4][R4.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0140*/ @!P0 LDG.E R10, desc[UR4][R6.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 84\n /*0150*/ LOP3.LUT P0, RZ, R12, 0x78, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 81\n /*0160*/ IMAD.WIDE R8, R11, 0x4, R8 ;\n /*0170*/ ISETP.LT.AND P0, PT, R11, UR6, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 365\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0180*/ ULDC UR6, c[0x0][0x234] ;\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*0190*/ FMUL R13, R13, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*01a0*/ FFMA R13, R13, UR6, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 87\n /*01b0*/ @!P0 EXIT ;\n /*01c0*/ STG.E desc[UR4][R8.64], R13 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 347\n\t//## File \".nv_debug_ptx_txt\", line 91\n /*01d0*/ EXIT ;\n.L_x_0:\n /*01e0*/ BRA `(.L_x_0);\n /*01f0*/ NOP;\n /*0200*/ NOP;\n /*0210*/ NOP;\n /*0220*/ NOP;\n /*0230*/ NOP;\n /*0240*/ NOP;\n /*0250*/ NOP;\n /*0260*/ NOP;\n /*0270*/ NOP;\n.L_x_1:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":346,"end_line":369,"code":" @triton.jit\n def fused_op_kernel(\n a_ptr,\n b_ptr,\n c_ptr,\n output_ptr,\n n_elements,\n scale_factor: float,\n ACTIVATION: tl.constexpr,\n BLOCK_SIZE: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n mask = offsets < n_elements\n\n a = tl.load(a_ptr + offsets, mask=mask)\n b = tl.load(b_ptr + offsets, mask=mask)\n c = tl.load(c_ptr + offsets, mask=mask)\n\n result = a * b * scale_factor + c\n if ACTIVATION == \"relu\":\n result = tl.where(result > 0, result, 0.0)\n\n tl.store(output_ptr + offsets, result, mask=mask)\n"}}} +{"event_type":"compilation","pid":4104992,"stack":[{"line":203,"name":"_run_module_as_main","filename":"","loc":""},{"line":88,"name":"_run_code","filename":"","loc":""},{"line":795,"name":"","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"raise SystemExit(main())"},{"line":783,"name":"main","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"artifacts = generate(workload, args.out_dir.resolve())"},{"line":655,"name":"generate","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"workload.run()"},{"line":431,"name":"_run_complex_kernels","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op(x, y, z, scale_factor=1.0, activation=\"relu\").sum()"},{"line":404,"name":"fused_op","filename":"/tritonparse/tritonparse/tools/generate_examples.py","loc":"fused_op_kernel[grid]("},{"line":374,"name":"","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"return lambda *args, **kwargs: self.run(grid=grid, warmup=False, *args, **kwargs)"},{"line":757,"name":"run","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self._do_compile(key, signature, device, constexprs, options, attrs, warmup)"},{"line":902,"name":"_do_compile","filename":"/python/lib/python3.14/site-packages/triton/runtime/jit.py","loc":"kernel = self.compile(src, target=target, options=options.__dict__)"},{"line":272,"name":"compile","filename":"/python/lib/python3.14/site-packages/triton/compiler/compiler.py","loc":"compilation_listener("},{"line":1640,"name":"maybe_trace_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"trace_structured_triton("},{"line":1554,"name":"trace_structured_triton","filename":"/tritonparse/tritonparse/structured_logging.py","loc":"metadata_dict[\"stack\"] = get_stack_trace()"}],"timestamp":"2026-09-23T10:04:05.%fZ","payload":{"metadata":{"cache_hit":true,"allowed_dot_input_precisions":["tf32","tf32x3","ieee","bf16x3","bf16x6"],"arch":"sm90","backend_name":"cuda","debug":false,"default_dot_input_precision":"tf32","deprecated_fp8_dot_operand_dtypes":["fp8e4b15"],"enable_fp_fusion":true,"enable_reflect_ftz":true,"extern_libs":[["libdevice","/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc"]],"global_scratch_align":1,"global_scratch_size":0,"hash":"cdc015d6a276c85a8b6f9cba0503bf70fe44d6f37b3877bf50efa3fe05c91037","instrumentation_mode":"","ir_override":null,"launch_cooperative_grid":false,"launch_pdl":false,"max_num_imprecise_acc_default":1073741824,"maxnreg":null,"name":"fused_op_kernel","num_ctas":1,"num_stages":3,"num_warps":4,"profile_scratch_align":1,"profile_scratch_size":0,"ptx_options":null,"ptx_version":null,"sanitize_overflow":true,"shared":0,"supported_fp8_dtypes":["fp8e4b15","fp8e4nv","fp8e5"],"target":{"backend":"cuda","arch":90,"warp_size":32},"tensordesc_meta":[],"tmem_size":0,"triton_version":"3.8.0","warp_size":32,"env":{},"src_attrs":{"(0,)":[["tt.divisibility",16]],"(1,)":[["tt.divisibility",16]],"(2,)":[["tt.divisibility",16]],"(3,)":[["tt.divisibility",16]],"(4,)":[],"(6,)":[],"(7,)":[]},"src_cache_key":"f2e5ae6c41d3bfc26af60eada8fff38a0b318cffe77df20b32d117346110ff07","src_constants":{"(6,)":"relu","(7,)":8},"times":{"ir_initialization":961,"lowering_stages":[],"store_results":0}},"file_path":{"fused_op_kernel.source":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.source","fused_op_kernel.ttir":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.ttir","fused_op_kernel.ttgir":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.ttgir","fused_op_kernel.llir":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.llir","fused_op_kernel.ptx":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.ptx","fused_op_kernel.cubin":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.cubin","fused_op_kernel.json":"/home/user/.triton/cache/ZXABLVVCO3EFVC3PTS5AKA57OD7EJVXTPM4HPP2Q56R74BOJCA3Q/fused_op_kernel.json"},"file_content":{"fused_op_kernel.ttir":"#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:5)\n#loc17 = loc(\"a_ptr\"(#loc))\n#loc18 = loc(\"b_ptr\"(#loc))\n#loc19 = loc(\"c_ptr\"(#loc))\n#loc20 = loc(\"output_ptr\"(#loc))\n#loc21 = loc(\"n_elements\"(#loc))\n#loc22 = loc(\"scale_factor\"(#loc))\nmodule {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %cst = arith.constant dense<0.000000e+00> : tensor<8xf32> loc(#loc1)\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc23)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc24)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32> loc(#loc25)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32> loc(#loc24)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32> loc(#loc24)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32> loc(#loc26)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32> loc(#loc26)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc27)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc27)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr> loc(#loc28)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc29)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc29)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr> loc(#loc30)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc31)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc31)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr> loc(#loc32)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32> loc(#loc33)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32> loc(#loc33)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32> loc(#loc33)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32> loc(#loc33)\n %result_13 = arith.cmpf ogt, %result_12, %cst : tensor<8xf32> loc(#loc34)\n %result_14 = arith.select %result_13, %result_12, %cst : tensor<8xi1>, tensor<8xf32> loc(#loc35)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr> loc(#loc15)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr>, tensor<8xi32> loc(#loc15)\n tt.store %1, %result_14, %mask_3 : tensor<8x!tt.ptr> loc(#loc16)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:15)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:19)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:38)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:16)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:13)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:21)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:13)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:21)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:13)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:18)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:31)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:22)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:18)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:9)\n#loc23 = loc(\"pid\"(#loc2))\n#loc24 = loc(\"offsets\"(#loc3))\n#loc25 = loc(\"offsets\"(#loc4))\n#loc26 = loc(\"mask\"(#loc5))\n#loc27 = loc(\"a\"(#loc6))\n#loc28 = loc(\"a\"(#loc7))\n#loc29 = loc(\"b\"(#loc8))\n#loc30 = loc(\"b\"(#loc9))\n#loc31 = loc(\"c\"(#loc10))\n#loc32 = loc(\"c\"(#loc11))\n#loc33 = loc(\"result\"(#loc12))\n#loc34 = loc(\"result\"(#loc13))\n#loc35 = loc(\"result\"(#loc14))\n","fused_op_kernel.ttgir":"#blocked = #ttg.blocked<{sizePerThread = [1], threadsPerWarp = [32], warpsPerCTA = [4], order = [0]}>\n#loc = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":347:5)\n#loc17 = loc(\"a_ptr\"(#loc))\n#loc18 = loc(\"b_ptr\"(#loc))\n#loc19 = loc(\"c_ptr\"(#loc))\n#loc20 = loc(\"output_ptr\"(#loc))\n#loc21 = loc(\"n_elements\"(#loc))\n#loc22 = loc(\"scale_factor\"(#loc))\nmodule attributes {\"ttg.num-ctas\" = 1 : i32, \"ttg.num-warps\" = 4 : i32, ttg.target = \"cuda:90\", \"ttg.threads-per-warp\" = 32 : i32} {\n tt.func public @fused_op_kernel(%a_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"a_ptr\"(#loc)), %b_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"b_ptr\"(#loc)), %c_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"c_ptr\"(#loc)), %output_ptr: !tt.ptr {tt.divisibility = 16 : i32} loc(\"output_ptr\"(#loc)), %n_elements: i32 loc(\"n_elements\"(#loc)), %scale_factor: f32 loc(\"scale_factor\"(#loc))) attributes {noinline = false} {\n %c8_i32 = arith.constant 8 : i32 loc(#loc1)\n %cst = arith.constant dense<0.000000e+00> : tensor<8xf32, #blocked> loc(#loc1)\n %pid = tt.get_program_id x : i32 loc(#loc23)\n %offsets = arith.muli %pid, %c8_i32 : i32 loc(#loc24)\n %offsets_0 = tt.make_range {end = 8 : i32, start = 0 : i32} : tensor<8xi32, #blocked> loc(#loc25)\n %offsets_1 = tt.splat %offsets : i32 -> tensor<8xi32, #blocked> loc(#loc24)\n %offsets_2 = arith.addi %offsets_1, %offsets_0 : tensor<8xi32, #blocked> loc(#loc24)\n %mask = tt.splat %n_elements : i32 -> tensor<8xi32, #blocked> loc(#loc26)\n %mask_3 = arith.cmpi slt, %offsets_2, %mask : tensor<8xi32, #blocked> loc(#loc26)\n %a = tt.splat %a_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc27)\n %a_4 = tt.addptr %a, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc27)\n %a_5 = tt.load %a_4, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc28)\n %b = tt.splat %b_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc29)\n %b_6 = tt.addptr %b, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc29)\n %b_7 = tt.load %b_6, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc30)\n %c = tt.splat %c_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc31)\n %c_8 = tt.addptr %c, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc31)\n %c_9 = tt.load %c_8, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc32)\n %result = arith.mulf %a_5, %b_7 : tensor<8xf32, #blocked> loc(#loc33)\n %result_10 = tt.splat %scale_factor : f32 -> tensor<8xf32, #blocked> loc(#loc33)\n %result_11 = arith.mulf %result, %result_10 : tensor<8xf32, #blocked> loc(#loc33)\n %result_12 = arith.addf %result_11, %c_9 : tensor<8xf32, #blocked> loc(#loc33)\n %result_13 = arith.cmpf ogt, %result_12, %cst : tensor<8xf32, #blocked> loc(#loc34)\n %result_14 = arith.select %result_13, %result_12, %cst : tensor<8xi1, #blocked>, tensor<8xf32, #blocked> loc(#loc35)\n %0 = tt.splat %output_ptr : !tt.ptr -> tensor<8x!tt.ptr, #blocked> loc(#loc15)\n %1 = tt.addptr %0, %offsets_2 : tensor<8x!tt.ptr, #blocked>, tensor<8xi32, #blocked> loc(#loc15)\n tt.store %1, %result_14, %mask_3 : tensor<8x!tt.ptr, #blocked> loc(#loc16)\n tt.return loc(#loc)\n } loc(#loc)\n} loc(#loc)\n#loc1 = loc(unknown)\n#loc2 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":357:15)\n#loc3 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:19)\n#loc4 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":358:38)\n#loc5 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":359:16)\n#loc6 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:21)\n#loc7 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":361:13)\n#loc8 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:21)\n#loc9 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":362:13)\n#loc10 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:21)\n#loc11 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":363:13)\n#loc12 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":365:18)\n#loc13 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:31)\n#loc14 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":367:22)\n#loc15 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:18)\n#loc16 = loc(\"/tritonparse/tritonparse/tools/generate_examples.py\":369:9)\n#loc23 = loc(\"pid\"(#loc2))\n#loc24 = loc(\"offsets\"(#loc3))\n#loc25 = loc(\"offsets\"(#loc4))\n#loc26 = loc(\"mask\"(#loc5))\n#loc27 = loc(\"a\"(#loc6))\n#loc28 = loc(\"a\"(#loc7))\n#loc29 = loc(\"b\"(#loc8))\n#loc30 = loc(\"b\"(#loc9))\n#loc31 = loc(\"c\"(#loc10))\n#loc32 = loc(\"c\"(#loc11))\n#loc33 = loc(\"result\"(#loc12))\n#loc34 = loc(\"result\"(#loc13))\n#loc35 = loc(\"result\"(#loc14))\n","fused_op_kernel.llir":"; ModuleID = 'LLVMDialectModule'\nsource_filename = \"LLVMDialectModule\"\ntarget datalayout = \"e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64\"\n\n; Function Attrs: nounwind\ndefine ptx_kernel void @fused_op_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, ptr addrspace(1) %2, ptr addrspace(1) %3, i32 %4, float %5, ptr addrspace(1) readnone captures(none) %6, ptr addrspace(1) readnone captures(none) %7) local_unnamed_addr #0 !dbg !4 {\n %9 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12\n %10 = shl i32 %9, 3, !dbg !13\n %11 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !14\n %12 = and i32 %11, 7, !dbg !14\n %13 = or disjoint i32 %10, %12, !dbg !13\n %14 = icmp slt i32 %13, %4, !dbg !15\n %15 = sext i32 %13 to i64, !dbg !16\n %16 = getelementptr [4 x i8], ptr addrspace(1) %0, i64 %15, !dbg !16\n %17 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %16, i1 %14) #2, !dbg !17\n %18 = bitcast i32 %17 to float, !dbg !17\n %19 = getelementptr [4 x i8], ptr addrspace(1) %1, i64 %15, !dbg !18\n %20 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %19, i1 %14) #2, !dbg !19\n %21 = bitcast i32 %20 to float, !dbg !19\n %22 = getelementptr [4 x i8], ptr addrspace(1) %2, i64 %15, !dbg !20\n %23 = tail call i32 asm sideeffect \"mov.u32 $0, 0x0;\\0A\\09@$2 ld.global.b32 { $0 }, [ $1 + 0 ];\", \"=r,l,b\"(ptr addrspace(1) %22, i1 %14) #2, !dbg !21\n %24 = bitcast i32 %23 to float, !dbg !21\n %25 = fmul float %18, %21, !dbg !22\n %26 = fmul float %5, %25, !dbg !22\n %27 = fadd float %26, %24, !dbg !22\n %28 = fcmp ogt float %27, 0.000000e+00, !dbg !23\n %29 = getelementptr [4 x i8], ptr addrspace(1) %3, i64 %15, !dbg !24\n %30 = and i32 %11, 120, !dbg !25\n %31 = icmp eq i32 %30, 0, !dbg !25\n %32 = bitcast float %27 to i32, !dbg !25\n %33 = select i1 %28, i32 %32, i32 0, !dbg !26\n %34 = and i1 %31, %14, !dbg !25\n tail call void asm sideeffect \"@$2 st.global.b32 [ $1 + 0 ], { $0 };\", \"r,l,b\"(i32 %33, ptr addrspace(1) %29, i1 %34) #2, !dbg !25\n ret void, !dbg !27\n}\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1\n\n; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)\ndeclare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1\n\nattributes #0 = { nounwind \"nvvm.reqntid\"=\"128\" }\nattributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }\nattributes #2 = { nounwind }\n\n!llvm.dbg.cu = !{!0}\n!llvm.module.flags = !{!2, !3}\n\n!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: \"triton\", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)\n!1 = !DIFile(filename: \"generate_examples.py\", directory: \"/tritonparse/tritonparse/tools\")\n!2 = !{i32 2, !\"Debug Info Version\", i32 3}\n!3 = !{i32 4, !\"nvvm-reflect-ftz\", i32 1}\n!4 = distinct !DISubprogram(name: \"fused_op_kernel\", linkageName: \"fused_op_kernel\", scope: !1, file: !1, line: 347, type: !5, scopeLine: 347, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)\n!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)\n!6 = !{null, !7, !7, !7, !7, !9, !8, !10, !10}\n!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !8, size: 64, dwarfAddressSpace: 1)\n!8 = !DIBasicType(name: \"float\", size: 32, encoding: DW_ATE_float)\n!9 = !DIBasicType(name: \"int\", size: 32, encoding: DW_ATE_signed)\n!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: \"pointer\", baseType: !11, size: 64, dwarfAddressSpace: 1)\n!11 = !DIBasicType(name: \"unknown_type\", encoding: DW_ATE_signed)\n!12 = !DILocation(line: 357, column: 15, scope: !4)\n!13 = !DILocation(line: 358, column: 19, scope: !4)\n!14 = !DILocation(line: 358, column: 38, scope: !4)\n!15 = !DILocation(line: 359, column: 16, scope: !4)\n!16 = !DILocation(line: 361, column: 21, scope: !4)\n!17 = !DILocation(line: 361, column: 13, scope: !4)\n!18 = !DILocation(line: 362, column: 21, scope: !4)\n!19 = !DILocation(line: 362, column: 13, scope: !4)\n!20 = !DILocation(line: 363, column: 21, scope: !4)\n!21 = !DILocation(line: 363, column: 13, scope: !4)\n!22 = !DILocation(line: 365, column: 18, scope: !4)\n!23 = !DILocation(line: 367, column: 31, scope: !4)\n!24 = !DILocation(line: 369, column: 18, scope: !4)\n!25 = !DILocation(line: 369, column: 9, scope: !4)\n!26 = !DILocation(line: 367, column: 22, scope: !4)\n!27 = !DILocation(line: 347, column: 5, scope: !4)\n","fused_op_kernel.ptx":"//\n// Generated by LLVM NVPTX Back-End\n//\n\n.version 8.8\n.target sm_90a\n.address_size 64\n\n\t// .globl\tfused_op_kernel // -- Begin function fused_op_kernel\n // @fused_op_kernel\n.visible .entry fused_op_kernel(\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_0,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_1,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_2,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_3,\n\t.param .u32 fused_op_kernel_param_4,\n\t.param .f32 fused_op_kernel_param_5,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_6,\n\t.param .u64 .ptr .global .align 1 fused_op_kernel_param_7\n)\n.reqntid 128\n{\n\t.reg .pred \t%p<5>;\n\t.reg .b32 \t%r<15>;\n\t.reg .b64 \t%rd<10>;\n\t.loc\t1 347 0 // generate_examples.py:347:0\n$L__func_begin0:\n\t.loc\t1 347 0 // generate_examples.py:347:0\n\n// %bb.0:\n\tld.param.b64 \t%rd5, [fused_op_kernel_param_0];\n\tld.param.b64 \t%rd6, [fused_op_kernel_param_1];\n$L__tmp0:\n\t.loc\t1 357 15 // generate_examples.py:357:15\n\tmov.u32 \t%r5, %ctaid.x;\n\t.loc\t1 358 19 // generate_examples.py:358:19\n\tshl.b32 \t%r6, %r5, 3;\n\tld.param.b64 \t%rd7, [fused_op_kernel_param_2];\n\tld.param.b64 \t%rd8, [fused_op_kernel_param_3];\n\t.loc\t1 358 38 // generate_examples.py:358:38\n\tmov.u32 \t%r7, %tid.x;\n\tand.b32 \t%r8, %r7, 7;\n\tld.param.b32 \t%r9, [fused_op_kernel_param_4];\n\t.loc\t1 358 19 // generate_examples.py:358:19\n\tor.b32 \t%r10, %r6, %r8;\n\tld.param.b32 \t%r11, [fused_op_kernel_param_5];\n\t.loc\t1 359 16 // generate_examples.py:359:16\n\tsetp.lt.s32 \t%p1, %r10, %r9;\n\t.loc\t1 361 21 // generate_examples.py:361:21\n\tmul.wide.s32 \t%rd9, %r10, 4;\n\tadd.s64 \t%rd1, %rd5, %rd9;\n\t.loc\t1 361 13 // generate_examples.py:361:13\n\t// begin inline asm\n\t.loc\t1 361 13 // generate_examples.py:361:13\n\tmov.u32 %r1, 0x0;\n\t@%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];\n\n\t// end inline asm\n\t.loc\t1 362 21 // generate_examples.py:362:21\n\tadd.s64 \t%rd2, %rd6, %rd9;\n\t.loc\t1 362 13 // generate_examples.py:362:13\n\t// begin inline asm\n\t.loc\t1 362 13 // generate_examples.py:362:13\n\tmov.u32 %r2, 0x0;\n\t@%p1 ld.global.b32 { %r2 }, [ %rd2 + 0 ];\n\n\t// end inline asm\n\t.loc\t1 363 21 // generate_examples.py:363:21\n\tadd.s64 \t%rd3, %rd7, %rd9;\n\t.loc\t1 363 13 // generate_examples.py:363:13\n\t// begin inline asm\n\t.loc\t1 363 13 // generate_examples.py:363:13\n\tmov.u32 %r3, 0x0;\n\t@%p1 ld.global.b32 { %r3 }, [ %rd3 + 0 ];\n\n\t// end inline asm\n\t.loc\t1 365 18 // generate_examples.py:365:18\n\tmul.f32 \t%r12, %r1, %r2;\n\tfma.rn.f32 \t%r13, %r11, %r12, %r3;\n\t.loc\t1 367 31 // generate_examples.py:367:31\n\tsetp.gt.f32 \t%p3, %r13, 0f00000000;\n\t.loc\t1 369 18 // generate_examples.py:369:18\n\tadd.s64 \t%rd4, %rd8, %rd9;\n\t.loc\t1 369 9 // generate_examples.py:369:9\n\tand.b32 \t%r14, %r7, 120;\n\tsetp.eq.b32 \t%p4, %r14, 0;\n\t.loc\t1 367 22 // generate_examples.py:367:22\n\tselp.b32 \t%r4, %r13, 0, %p3;\n\t.loc\t1 369 9 // generate_examples.py:369:9\n\tand.pred \t%p2, %p4, %p1;\n\t// begin inline asm\n\t@%p2 st.global.b32 [ %rd4 + 0 ], { %r4 };\n\n\t// end inline asm\n\t.loc\t1 347 5 // generate_examples.py:347:5\n\tret;\n$L__tmp1:\n$L__func_end0:\n // -- End function\n}\n\t.file\t1 \"/tritonparse/tritonparse/tools/generate_examples.py\"\n\t.section\t.debug_abbrev\n\t{\n.b8 1 // Abbreviation Code\n.b8 17 // DW_TAG_compile_unit\n.b8 0 // DW_CHILDREN_no\n.b8 37 // DW_AT_producer\n.b8 8 // DW_FORM_string\n.b8 19 // DW_AT_language\n.b8 5 // DW_FORM_data2\n.b8 3 // DW_AT_name\n.b8 8 // DW_FORM_string\n.b8 16 // DW_AT_stmt_list\n.b8 6 // DW_FORM_data4\n.b8 27 // DW_AT_comp_dir\n.b8 8 // DW_FORM_string\n.b8 0 // EOM(1)\n.b8 0 // EOM(2)\n.b8 0 // EOM(3)\n\t}\n\t.section\t.debug_info\n\t{\n.b32 113 // Length of Unit\n.b8 2 // DWARF version number\n.b8 0\n.b32 .debug_abbrev // Offset Into Abbrev. Section\n.b8 8 // Address Size (in bytes)\n.b8 1 // Abbrev [1] 0xb:0x6a DW_TAG_compile_unit\n.b8 116 // DW_AT_producer\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 0\n.b8 2 // DW_AT_language\n.b8 0\n.b8 103 // DW_AT_name\n.b8 101\n.b8 110\n.b8 101\n.b8 114\n.b8 97\n.b8 116\n.b8 101\n.b8 95\n.b8 101\n.b8 120\n.b8 97\n.b8 109\n.b8 112\n.b8 108\n.b8 101\n.b8 115\n.b8 46\n.b8 112\n.b8 121\n.b8 0\n.b32 .debug_line // DW_AT_stmt_list\n.b8 47 // DW_AT_comp_dir\n.b8 100\n.b8 97\n.b8 116\n.b8 97\n.b8 47\n.b8 117\n.b8 115\n.b8 101\n.b8 114\n.b8 115\n.b8 47\n.b8 121\n.b8 104\n.b8 97\n.b8 111\n.b8 47\n.b8 102\n.b8 98\n.b8 115\n.b8 111\n.b8 117\n.b8 114\n.b8 99\n.b8 101\n.b8 47\n.b8 102\n.b8 98\n.b8 99\n.b8 111\n.b8 100\n.b8 101\n.b8 47\n.b8 112\n.b8 121\n.b8 116\n.b8 111\n.b8 114\n.b8 99\n.b8 104\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 114\n.b8 105\n.b8 116\n.b8 111\n.b8 110\n.b8 112\n.b8 97\n.b8 114\n.b8 115\n.b8 101\n.b8 47\n.b8 116\n.b8 111\n.b8 111\n.b8 108\n.b8 115\n.b8 0\n\t}\n\t.section\t.debug_macinfo\t{\t}\n","fused_op_kernel.json":"{\"hash\": \"cdc015d6a276c85a8b6f9cba0503bf70fe44d6f37b3877bf50efa3fe05c91037\", \"target\": {\"backend\": \"cuda\", \"arch\": 90, \"warp_size\": 32}, \"num_warps\": 4, \"num_ctas\": 1, \"num_stages\": 3, \"warp_size\": 32, \"maxnreg\": null, \"ptx_version\": null, \"ptx_options\": null, \"ir_override\": null, \"enable_fp_fusion\": true, \"enable_reflect_ftz\": true, \"launch_cooperative_grid\": false, \"launch_pdl\": false, \"supported_fp8_dtypes\": [\"fp8e4b15\", \"fp8e4nv\", \"fp8e5\"], \"deprecated_fp8_dot_operand_dtypes\": [\"fp8e4b15\"], \"default_dot_input_precision\": \"tf32\", \"allowed_dot_input_precisions\": [\"tf32\", \"tf32x3\", \"ieee\", \"bf16x3\", \"bf16x6\"], \"max_num_imprecise_acc_default\": 1073741824, \"extern_libs\": [[\"libdevice\", \"/python/lib/python3.14/site-packages/triton/backends/nvidia/lib/libdevice.10.bc\"]], \"debug\": false, \"backend_name\": \"cuda\", \"sanitize_overflow\": true, \"arch\": \"sm90\", \"instrumentation_mode\": \"\", \"triton_version\": \"3.8.0\", \"tensordesc_meta\": [], \"shared\": 0, \"tmem_size\": 0, \"global_scratch_size\": 0, \"global_scratch_align\": 1, \"profile_scratch_size\": 0, \"profile_scratch_align\": 1, \"name\": \"fused_op_kernel\"}","fused_op_kernel.sass":"Function:fused_op_kernel\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 347\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0000*/ LDC R1, c[0x0][0x28] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 41\n /*0010*/ S2R R12, SR_TID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*0020*/ LDC.64 R2, c[0x0][0x210] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*0030*/ ULDC UR6, c[0x0][0x230] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0040*/ MOV R13, RZ ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 357\n\t//## File \".nv_debug_ptx_txt\", line 35\n /*0050*/ S2R R11, SR_CTAID.X ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 60\n /*0060*/ LDC.64 R4, c[0x0][0x218] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0070*/ LDC.64 R6, c[0x0][0x220] ;\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0080*/ HFMA2.MMA R10, -RZ, RZ, 0, 0 ;\n\t//## File \".nv_debug_ptx_txt\", line 22\n /*0090*/ ULDC.64 UR4, c[0x0][0x208] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 83\n /*00a0*/ LDC.64 R8, c[0x0][0x228] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 358\n\t//## File \".nv_debug_ptx_txt\", line 42\n /*00b0*/ LOP3.LUT R0, R12, 0x7, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 45\n /*00c0*/ LEA R11, R11, R0, 0x3 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 56\n /*00d0*/ HFMA2.MMA R0, -RZ, RZ, 0, 0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 359\n\t//## File \".nv_debug_ptx_txt\", line 48\n /*00e0*/ ISETP.GE.AND P0, PT, R11.reuse, UR6, PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 51\n /*00f0*/ IMAD.WIDE R2, R11, 0x4, R2 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 60\n /*0100*/ IMAD.WIDE R4, R11, 0x4, R4 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 69\n /*0110*/ IMAD.WIDE R6, R11, 0x4, R6 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 361\n\t//## File \".nv_debug_ptx_txt\", line 56\n /*0120*/ @!P0 LDG.E R0, desc[UR4][R2.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 362\n\t//## File \".nv_debug_ptx_txt\", line 65\n /*0130*/ @!P0 LDG.E R13, desc[UR4][R4.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 363\n\t//## File \".nv_debug_ptx_txt\", line 74\n /*0140*/ @!P0 LDG.E R10, desc[UR4][R6.64] ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 86\n /*0150*/ LOP3.LUT P0, RZ, R12, 0x78, RZ, 0xc0, !PT ;\n\t//## File \".nv_debug_ptx_txt\", line 83\n /*0160*/ IMAD.WIDE R2, R11, 0x4, R8 ;\n /*0170*/ ISETP.LT.AND P0, PT, R11, UR6, !P0 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 365\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*0180*/ ULDC UR6, c[0x0][0x234] ;\n\t//## File \".nv_debug_ptx_txt\", line 78\n /*0190*/ FMUL R13, R13, R0 ;\n\t//## File \".nv_debug_ptx_txt\", line 79\n /*01a0*/ FFMA R10, R13, UR6, R10 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 367\n\t//## File \".nv_debug_ptx_txt\", line 88\n /*01b0*/ FMNMX R9, RZ, R10, !PT ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 369\n\t//## File \".nv_debug_ptx_txt\", line 92\n /*01c0*/ @!P0 EXIT ;\n /*01d0*/ STG.E desc[UR4][R2.64], R9 ;\n\t//## File \"/tritonparse/tritonparse/tools/generate_examples.py\", line 347\n\t//## File \".nv_debug_ptx_txt\", line 96\n /*01e0*/ EXIT ;\n.L_x_0:\n /*01f0*/ BRA `(.L_x_0);\n /*0200*/ NOP;\n /*0210*/ NOP;\n /*0220*/ NOP;\n /*0230*/ NOP;\n /*0240*/ NOP;\n /*0250*/ NOP;\n /*0260*/ NOP;\n /*0270*/ NOP;\n.L_x_1:\n\n\n//--------------------- SYMBOLS --------------------------\n\n\t.type\t\t.nv.reservedSmem.offset0,@object\n\t.size\t\t.nv.reservedSmem.offset0,0x4\n\n"},"python_source":{"file_path":"/tritonparse/tritonparse/tools/generate_examples.py","start_line":346,"end_line":369,"code":" @triton.jit\n def fused_op_kernel(\n a_ptr,\n b_ptr,\n c_ptr,\n output_ptr,\n n_elements,\n scale_factor: float,\n ACTIVATION: tl.constexpr,\n BLOCK_SIZE: tl.constexpr,\n ):\n pid = tl.program_id(axis=0)\n offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n mask = offsets < n_elements\n\n a = tl.load(a_ptr + offsets, mask=mask)\n b = tl.load(b_ptr + offsets, mask=mask)\n c = tl.load(c_ptr + offsets, mask=mask)\n\n result = a * b * scale_factor + c\n if ACTIVATION == \"relu\":\n result = tl.where(result > 0, result, 0.0)\n\n tl.store(output_ptr + offsets, result, mask=mask)\n"}}} diff --git a/tests/example_output/parsed_output_complex/dedicated_log_triton_trace_findhao__mapped.ndjson.gz b/tests/example_output/parsed_output_complex/dedicated_log_triton_trace_findhao__mapped.ndjson.gz index 81514a4..833dc93 100644 Binary files a/tests/example_output/parsed_output_complex/dedicated_log_triton_trace_findhao__mapped.ndjson.gz and b/tests/example_output/parsed_output_complex/dedicated_log_triton_trace_findhao__mapped.ndjson.gz differ diff --git a/tests/example_output/parsed_output_inductor/dedicated_log_triton_trace_inductor__mapped.ndjson.gz b/tests/example_output/parsed_output_inductor/dedicated_log_triton_trace_inductor__mapped.ndjson.gz index 547a175..4dfb43e 100644 Binary files a/tests/example_output/parsed_output_inductor/dedicated_log_triton_trace_inductor__mapped.ndjson.gz and b/tests/example_output/parsed_output_inductor/dedicated_log_triton_trace_inductor__mapped.ndjson.gz differ diff --git a/tritonparse/parse/event_diff.py b/tritonparse/parse/event_diff.py index 07ab249..f8f730e 100644 --- a/tritonparse/parse/event_diff.py +++ b/tritonparse/parse/event_diff.py @@ -1,7 +1,10 @@ # Copyright (c) Meta Platforms, Inc. and affiliates. +import ast +import hashlib +import re from collections import defaultdict, OrderedDict -from typing import Any, Dict, List, Optional, Tuple +from typing import Any, Dict, List, Optional, Set, Tuple from tritonparse._json_compat import dumps, loads @@ -108,9 +111,715 @@ def _dedup_compilations_by_hash( return deduped +def _is_tensor_arg_value(arg_val: Any) -> bool: + """True when an extracted arg value describes a plain tensor argument.""" + return isinstance(arg_val, dict) and arg_val.get("type") == "tensor" + + +def _is_tensor_like_value(arg_val: Any) -> bool: + """True for tensor args including descriptors and foreign wrappers. + + Covers plain tensors plus TMA TensorDescriptors and triton_kernels + Tensor/Storage wrappers, which nest backing-tensor metadata. + """ + if not isinstance(arg_val, dict): + return False + arg_type = arg_val.get("type", "") + return ( + arg_type == "tensor" + or arg_type == "TensorDescriptor" + or (isinstance(arg_type, str) and arg_type.startswith("triton_kernels.tensor.")) + ) + + +def _extract_scalar_value(arg_val: Any) -> Any: + """Unwrap {"type": ..., "value": ...} scalar wrappers, if present. + + Falls back to the "repr" payload for values recorded without a + structured value (None, tuples, other objects), so config values + such as "None" still compare against best_config spellings. + """ + if isinstance(arg_val, dict) and "type" in arg_val and "value" in arg_val: + return arg_val["value"] + if isinstance(arg_val, dict) and "repr" in arg_val: + return arg_val["repr"] + return arg_val + + +def _is_note_only_launch(extracted: Dict[str, Any]) -> bool: + """True when a launch carries no extractable arguments. + + During CUDA graph capture, argument extraction is skipped and the + launch records {"_note": ...} instead of real args. Such launches + carry no key information (a plain string marker, never a wrapped + {"type", "value"} kernel parameter), so key partitioning ignores + them; the same round's warmup launches outside capture still carry + the full arguments. + """ + return isinstance(extracted.get("_note"), str) + + +def _drop_data_ptrs(value: Any) -> Any: + """Recursively drop data_ptr fields from extracted argument values. + + Tensor descriptors and foreign tensor wrappers nest backing-tensor + metadata (including volatile data_ptrs) inside the argument dict. + Dropping data_ptrs keeps signatures stable across reallocations + while retaining identity metadata (shapes, dtypes, block shapes). + """ + if isinstance(value, dict): + return {k: _drop_data_ptrs(v) for k, v in value.items() if k != "data_ptr"} + if isinstance(value, list): + return [_drop_data_ptrs(v) for v in value] + return value + + +def _fingerprint_value(value: Any) -> str: + """Stable fingerprint of one config-detection leaf value.""" + try: + return dumps(_drop_data_ptrs(value), sort_keys=True) + except TypeError: + return dumps(str(value)) + + +def _config_leaf_fingerprints(arg_val: Any) -> List[Tuple[Tuple[str, ...], str]]: + """Locate config correlation at (path, fingerprint) granularity. + + Nested dicts (descriptors, foreign wrappers) are fingerprinted per + leaf path so a config-shaped nested field (e.g. block_shape) does + not drag key-shaped siblings (e.g. shape) out of the signature. + Lists are leaves: per-index splits are contrived, and whole-list + comparison already matches the legacy behavior when a list mixes + key/config content. data_ptr entries are skipped everywhere since + they are volatile across reallocations. + """ + if not isinstance(arg_val, dict): + return [((), _fingerprint_value(arg_val))] + flattened: List[Tuple[Tuple[str, ...], str]] = [] + stack: List[Tuple[Dict[str, Any], Tuple[str, ...]]] = [(arg_val, ())] + while stack: + current, path = stack.pop() + for key in sorted(current, key=str): + if key == "data_ptr": + continue + child = current[key] + child_path = path + (key,) + if isinstance(child, dict): + stack.append((child, child_path)) + else: + flattened.append((child_path, _fingerprint_value(child))) + return flattened + + +def _mask_config_paths( + value: Dict[str, Any], config_paths: Set[Tuple[str, ...]] +) -> Dict[str, Any]: + """Copy a nested arg value with config-correlated leaves removed. + + data_ptr fields are dropped first (see _drop_data_ptrs); paths that + are absent in this launch's shape are skipped. + """ + masked = _drop_data_ptrs(value) + for path in config_paths: + current = masked + for step in path[:-1]: + if not isinstance(current, dict) or step not in current: + current = None + break + current = current[step] + if path and isinstance(current, dict): + current.pop(path[-1], None) + return masked + + +# Marks a nested path missing from a launch's arg shape during config +# detection. Presence/absence correlated with compilation is itself a +# config signal (e.g. "value" vs "repr" in None/int scalar wrappers). +_ABSENT: Any = object() + + +def _find_config_args( + group_infos: List[Tuple[str, Optional[str], Dict[str, Any]]], +) -> Dict[str, Optional[Set[Tuple[str, ...]]]]: + """Find launch args (or nested fields) determined by compilation hash. + + A (arg, path) pair is config-correlated when every compilation hash + in the session shows exactly one fingerprint for it, while distinct + hashes disagree (e.g. BLOCK_SIZE_M). Tensor args never participate. + Launches without a compilation hash cannot be attributed, so they + are skipped here. + + Returns a mapping of arg name to the config-correlated paths within + it, or None when the whole arg is config-correlated (scalars and + fully config-shaped descriptors). + """ + values_by_path: Dict[Tuple[str, Tuple[str, ...]], Dict[str, Set[Any]]] = ( + defaultdict(lambda: defaultdict(set)) + ) + paths_by_arg: Dict[str, Set[Tuple[str, ...]]] = defaultdict(set) + hash_universe: Set[str] = set() + for _group_hash, comp_hash, extracted in group_infos: + if not comp_hash: + continue + hash_universe.add(comp_hash) + for arg_name, arg_val in extracted.items(): + # Plain tensors never participate: tensor identity is key-like + # (shapes group rounds), never config-like. Descriptors and + # foreign wrappers DO participate on data_ptr-free content, so + # config-shaped fields (e.g. a pre-hook rewriting block_shape + # per config) are still recognized as config params. + if _is_tensor_arg_value(arg_val): + continue + for path, fingerprint in _config_leaf_fingerprints(arg_val): + paths_by_arg[arg_name].add(path) + values_by_path[(arg_name, path)][comp_hash].add(fingerprint) + for (_arg_name, path), per_hash in values_by_path.items(): + if not path: + continue + # A nested path missing under a hash counts as a distinct value; + # otherwise a field present under only one hash would escape + # config detection and split the session by config. + for comp_hash in hash_universe: + if comp_hash not in per_hash: + per_hash[comp_hash] = {_ABSENT} + config_paths_by_arg: Dict[str, Set[Tuple[str, ...]]] = defaultdict(set) + for (arg_name, path), per_hash in values_by_path.items(): + if len(per_hash) <= 1: + continue + if not all(len(values) == 1 for values in per_hash.values()): + continue + # The values must actually disagree across hashes; a constant + # field is not a config param even though it trivially has "one + # value per hash". + if len(set().union(*per_hash.values())) > 1: + config_paths_by_arg[arg_name].add(path) + config_args: Dict[str, Optional[Set[Tuple[str, ...]]]] = {} + for arg_name, paths in config_paths_by_arg.items(): + if paths and paths == paths_by_arg[arg_name]: + config_args[arg_name] = None + else: + config_args[arg_name] = paths + return config_args + + +def _launch_key_signature( + extracted: Dict[str, Any], + config_args: Dict[str, Optional[Set[Tuple[str, ...]]]], +) -> str: + """Stable signature of a launch's autotune-key arguments. + + Covers all scalar args except autotune config params, plus tensor + identity metadata (dtype/shape/strides; data_ptr excluded since it + changes on every allocation). Launches from one key invocation share a + signature; different keys differ in at least one runtime arg. + """ + payload: Dict[str, Any] = {} + for arg_name in sorted(extracted): + arg_val = extracted[arg_name] + if arg_name in config_args: + paths = config_args[arg_name] + if paths is None or not isinstance(arg_val, dict): + continue + # Mixed key/config descriptor: mask only the config paths so + # key-shaped nested fields still separate key rounds. + payload[arg_name] = _mask_config_paths(arg_val, paths) + elif _is_tensor_arg_value(arg_val): + payload[arg_name] = { + key: arg_val.get(key) + for key in ("dtype", "shape", "stride", "strides", "numel") + if key in arg_val + } + elif isinstance(arg_val, dict): + # Descriptors and foreign wrappers: keep identity metadata + # but drop volatile nested data_ptrs for stability. + payload[arg_name] = _drop_data_ptrs(arg_val) + else: + payload[arg_name] = arg_val + try: + serialized = dumps(payload, sort_keys=True) + except TypeError: + serialized = dumps(str(payload)) + return hashlib.sha256(serialized.encode("utf-8")).hexdigest()[:16] + + +def _split_session_by_launch_key( + session_id: str, + session_data: Dict[str, Any], + launch_by_group_hash: Dict[str, Dict[str, Any]], +) -> List[Dict[str, Any]]: + """Partition one call-site session into per-autotune-key sub-sessions. + + Repeated ``Autotuner.run()`` invocations at the same call site (e.g. a + loop over problem sizes) share one call-site session id. Each invocation + benchmarks with identical runtime args, so grouping launches by key + signature recovers the original invocations. + + Single-partition sessions keep the plain coarse session id so their + output is unchanged; split sessions get "{session_id}:{full_sig}" + ids (the full 16-char signature: truncating it could collide and + silently drop a key round). Partitions are ordered by first + occurrence for stable output. Each partition carries + launch_group_hashes (ordered as the legacy path orders them) and + benchmark/winner occurrence ids. + + Limitation: launches without distinguishing arguments cannot be + attributed to a key (``extracted_args`` is optional in the launch + schema). Distinct listener keys alone cannot partition + indistinguishable launches, so such rounds share one partition and + only the latest listener result is attached. Likewise, an autotune + key on a tl.constexpr/specialized argument is misread as a config + param (one value per compilation hash) whenever tensor shapes do + not already separate the rounds; deriving config fields from + best_config is impossible on traces without listener events, and + compilation metadata does not record constexpr values. Note-only + launches in a split session stay unattributed: their occurrence ids + are excluded from every partition since no key can be derived for + them. Result matching is heuristic as well: a result matches when + every best_config pair resolvable against the winner launch agrees, + so rounds whose distinguishing config fields are all unresolvable + are told apart by shared fields and file order only. + """ + all_occurrences = session_data.get("launch_occurrences", []) or [] + group_hashes = session_data.get("launch_group_hashes", set()) or set() + + group_infos: List[Tuple[str, Optional[str], Dict[str, Any]]] = [] + skipped_note_groups = set() + for group_hash in group_hashes: + launch = launch_by_group_hash.get(group_hash, {}) + comp_hash = launch.get("compilation_metadata", {}).get("hash") + extracted = launch.get("extracted_args", {}) or {} + # Launches captured inside CUDA graphs carry a _note marker + # instead of real args. They hold no key information and would + # pool unrelated rounds into one partition, so key derivation + # ignores them; emission still counts their occurrences whenever + # no attribution choice is needed (see below). + if _is_note_only_launch(extracted): + skipped_note_groups.add(group_hash) + continue + group_infos.append((group_hash, comp_hash, extracted)) + occurrences = [ + record + for record in all_occurrences + if record.get("launch_group_hash") not in skipped_note_groups + ] + + def _single_partition( + benchmark_ids: List[int], winner_ids: List[int] + ) -> List[Dict[str, Any]]: + ordered = sorted( + (info[0] for info in group_infos), + key=lambda h: launch_by_group_hash.get(h, {}).get("occurrence_id", 0), + ) + return [ + { + "sub_session_id": session_id, + "launch_group_hashes": ordered, + "benchmark_occurrence_ids": benchmark_ids, + "winner_occurrence_ids": winner_ids, + } + ] + + if not group_infos: + # Nothing attributable to split on: either a compilations-only + # session (benchmark launches untraced) or an all-note session + # (every launch captured inside CUDA graphs). Preserve the coarse + # session verbatim (legacy parity): the groups keep output-file + # lookup and the neither-compilations-nor-launches guard working + # when the kernels compiled outside this call site, and the coarse + # occurrence ids keep the real launches instead of reporting an + # empty, cache-looking session. + ordered = sorted( + group_hashes, + key=lambda h: launch_by_group_hash.get(h, {}).get("occurrence_id", 0), + ) + return [ + { + "sub_session_id": session_id, + "launch_group_hashes": ordered, + "benchmark_occurrence_ids": list( + session_data.get("benchmark_occurrence_ids", []) + ), + "winner_occurrence_ids": list( + session_data.get("winner_occurrence_ids", []) + ), + } + ] + if not occurrences: + # Defensive: same-process ingest always records occurrences, but + # keep emission working (with coarse occurrence lists) if absent. + return _single_partition( + list(session_data.get("benchmark_occurrence_ids", [])), + list(session_data.get("winner_occurrence_ids", [])), + ) + + config_args = _find_config_args(group_infos) + signature_by_group = { + group_hash: _launch_key_signature(extracted, config_args) + for group_hash, _comp_hash, extracted in group_infos + } + grouped: Dict[str, List[str]] = defaultdict(list) + for group_hash, _comp_hash, _extracted in group_infos: + grouped[signature_by_group[group_hash]].append(group_hash) + + if len(grouped) == 1: + # One key round: no attribution choice, so note occurrences join + # the emission (their arg payload stays excluded from the groups + # above, keeping _note markers out of the varies table). + bench_ids = [ + record["occurrence_id"] + for record in all_occurrences + if record.get("is_benchmark") + ] + winner_ids = [ + record["occurrence_id"] + for record in all_occurrences + if not record.get("is_benchmark") + ] + return _single_partition(bench_ids, winner_ids) + + def _partition_min_occurrence(sig: str) -> int: + member_groups = set(grouped[sig]) + occs = [ + record["occurrence_id"] + for record in occurrences + if record.get("launch_group_hash") in member_groups + ] + return min(occs) if occs else 0 + + partitions = [] + for sig in sorted(grouped, key=_partition_min_occurrence): + member_groups = set(grouped[sig]) + bench_ids = [] + winner_ids = [] + for record in occurrences: + if record.get("launch_group_hash") not in member_groups: + continue + if record.get("is_benchmark"): + bench_ids.append(record["occurrence_id"]) + else: + winner_ids.append(record["occurrence_id"]) + ordered = sorted( + grouped[sig], + key=lambda h: launch_by_group_hash.get(h, {}).get("occurrence_id", 0), + ) + partitions.append( + { + "sub_session_id": f"{session_id}:{sig}", + "launch_group_hashes": ordered, + "benchmark_occurrence_ids": bench_ids, + "winner_occurrence_ids": winner_ids, + } + ) + return partitions + + +def _resolve_partition_winner( + winner_occurrence_ids: List[int], + occurrence_to_group: Dict[int, str], + launch_by_group_hash: Dict[str, Dict[str, Any]], +) -> Tuple[Optional[str], Optional[str]]: + """Resolve (selected launch_group_hash, winner compilation hash). + + Mirrors the legacy last-non-benchmark-launch rule, scoped to one + partition: the latest winner occurrence whose launch carries a hash. + """ + for occurrence_id in reversed(winner_occurrence_ids): + group_hash = occurrence_to_group.get(occurrence_id) + if not group_hash: + continue + launch = launch_by_group_hash.get(group_hash, {}) + comp_hash = launch.get("compilation_metadata", {}).get("hash") + if comp_hash: + return group_hash, comp_hash + return None, None + + +# A value runs until the next "name:" pair (or end of string), so tuple/list +# values containing commas are kept intact. +_BEST_CONFIG_PAIR_RE = re.compile( + r"([A-Za-z_]\w*)\s*:\s*(.*?)(?=,\s*[A-Za-z_]\w*\s*:|$)" +) + + +def _parse_best_config(best_config: Any) -> Dict[str, str]: + """Parse Triton Config str form ("BLOCK_SIZE_M: 16, num_warps: 1, ...").""" + if not isinstance(best_config, str): + return {} + return { + match.group(1): match.group(2).strip() + for match in _BEST_CONFIG_PAIR_RE.finditer(best_config) + } + + +def _parse_sequence_literal(text: str) -> Optional[List[Any]]: + """Parse a best_config sequence spelling ("(2, 1, 1)") into a list.""" + try: + value = ast.literal_eval(text) + except (ValueError, SyntaxError): + return None + if isinstance(value, (list, tuple)): + return list(value) + return None + + +def _config_value_matches(expected: str, actual: Any) -> bool: + """Compare a parsed best_config value against a launch value. + + Int-aware first (exact for large ints; also accepts hex), then + sequences (best_config prints tuples while JSON round-trips turn + them into lists), then numeric, then string fallback. Bools compare + against true/false/1/0 spellings. + """ + text = expected.strip() + if actual is None: + return text in ("None", "null", "") + if isinstance(actual, bool): + normalized = text.lower() + if normalized in ("1", "true"): + return actual is True + if normalized in ("0", "false"): + return actual is False + return False + if isinstance(actual, int): + try: + return int(text, 0) == actual + except ValueError: + pass + if isinstance(actual, (list, tuple)): + parsed = _parse_sequence_literal(text) + return parsed is not None and parsed == list(actual) + try: + return float(text) == float(actual) + except (TypeError, ValueError): + return text == str(actual) + + +def _match_autotune_result( + results: List[Dict[str, Any]], + winner_launch: Dict[str, Any], + skip_indices: Set[int], + winner_warps_base: Optional[int] = None, +) -> Optional[int]: + """Match listener results to a partition via its winner launch. + + Every best_config pair resolvable against the winner's config values + (extracted scalar args, then compilation metadata) must agree. + Tensor-like extracted args are skipped (a config name can never + denote a tensor). Under warp specialization the launch metadata + carries the expanded num_warps while best_config prints the + requested one, so num_warps compares against winner_warps_base when + the compilation payload recovered one. Returns the index of the + first match in file order that is not in skip_indices, so each + result is consumed by at most one partition. + + This is heuristic: pairs unresolvable against the winner (e.g. BLOCK + sizes, which are neither kernel args nor compilation metadata) are + skipped, so a match on shared fields alone can misattribute results + when file order disagrees with round order. + """ + extracted = winner_launch.get("extracted_args", {}) or {} + compilation_metadata = winner_launch.get("compilation_metadata", {}) or {} + for index, result in enumerate(results): + if index in skip_indices: + continue + pairs = _parse_best_config(result.get("best_config")) + compared = 0 + matched = 0 + for name, expected in pairs.items(): + if name == "num_warps" and winner_warps_base is not None: + actual = winner_warps_base + elif name in extracted and not _is_tensor_like_value(extracted[name]): + actual = _extract_scalar_value(extracted[name]) + elif name in compilation_metadata: + actual = compilation_metadata[name] + else: + continue + compared += 1 + if _config_value_matches(expected, actual): + matched += 1 + if compared > 0 and matched == compared: + return index + return None + + +def _globally_referenced_hashes( + partitions: List[Dict[str, Any]], + launch_by_group_hash: Dict[str, Dict[str, Any]], +) -> Set[str]: + """Compilation hashes referenced by any partition's launches.""" + referenced = set() + for part in partitions: + for group_hash in part["launch_group_hashes"]: + launch = launch_by_group_hash.get(group_hash, {}) + comp_hash = launch.get("compilation_metadata", {}).get("hash") + if comp_hash: + referenced.add(comp_hash) + return referenced + + +def _partition_compilations( + part: Dict[str, Any], + coarse_compilations: List[Dict[str, Any]], + launch_by_group_hash: Dict[str, Dict[str, Any]], + unattributed_hashes: Set[str], +) -> List[Dict[str, Any]]: + """Compilations for a partition, in coarse order. + + Besides compilations referenced by this partition's launches, this + keeps compilations unattributed to any partition (e.g. a candidate + whose benchmark launches were never traced), so a split never + silently drops a config that the unsplit session would have shown. + """ + referenced = set(unattributed_hashes) + for group_hash in part["launch_group_hashes"]: + launch = launch_by_group_hash.get(group_hash, {}) + comp_hash = launch.get("compilation_metadata", {}).get("hash") + if comp_hash: + referenced.add(comp_hash) + return [ + comp + for comp in coarse_compilations + if comp.get("payload", {}).get("metadata", {}).get("hash") in referenced + ] + + +def _build_sub_session( + part: Dict[str, Any], + coarse_compilations: List[Dict[str, Any]], + occurrence_to_group: Dict[int, str], + results: List[Dict[str, Any]], + used_result_indices: Set[int], + unattributed_hashes: Set[str], + single: bool, + launch_by_group_hash: Dict[str, Dict[str, Any]], +) -> Tuple[Dict[str, Any], Optional[str]]: + """Build one expanded sub-session entry plus its selected group hash.""" + if single: + # No split: keep every coarse compilation (including ones no + # launch references) so output matches the legacy path. + part_compilations = list(coarse_compilations) + else: + part_compilations = _partition_compilations( + part, coarse_compilations, launch_by_group_hash, unattributed_hashes + ) + + selected_group, winner_hash = _resolve_partition_winner( + part["winner_occurrence_ids"], + occurrence_to_group, + launch_by_group_hash, + ) + + matched_result: Optional[Dict[str, Any]] = None + if results: + if single: + # Legacy parity: ingest used to keep only the last result. + matched_result = results[-1] + elif selected_group: + # Warp-specialized kernels record the requested num_warps + # as num_warps_base on the compilation payload; the launch + # metadata only carries the expanded count. + winner_warps_base = None + if winner_hash: + for comp in part_compilations: + comp_meta = comp.get("payload", {}).get("metadata", {}) + if comp_meta.get("hash") == winner_hash: + winner_warps_base = comp_meta.get("num_warps_base") + break + match_index = _match_autotune_result( + results, + launch_by_group_hash.get(selected_group, {}), + used_result_indices, + winner_warps_base, + ) + if match_index is not None: + used_result_indices.add(match_index) + matched_result = results[match_index] + + sub_data: Dict[str, Any] = { + "compilations": part_compilations, + "launch_group_hashes": set(part["launch_group_hashes"]), + "benchmark_occurrence_ids": part["benchmark_occurrence_ids"], + "winner_occurrence_ids": part["winner_occurrence_ids"], + } + if matched_result: + sub_data["autotune_result"] = matched_result + return sub_data, selected_group + + +def _expand_sessions_by_launch_key( + autotune_sessions: Dict[str, Dict[str, Any]], + session_stacks: Dict[str, List[Dict[str, Any]]], + launch_by_group_hash: Dict[str, Dict[str, Any]], +) -> Tuple[Dict[str, Dict[str, Any]], Dict[str, str], Dict[str, List[Dict[str, Any]]]]: + """Expand call-site sessions into per-autotune-key sub-sessions. + + Returns (expanded_sessions, expanded_winners, expanded_stacks) with the + same shapes the emission passes already consume: each sub-session keeps + compilations, launch_group_hashes, and occurrence id lists; winners map + to the selected launch_group_hash; stacks still describe the shared + call site. Each sub-session also carries its own "autotune_result" + (matched by best config with each result consumed at most once, or + the last result when no split happened). + """ + expanded_sessions: Dict[str, Dict[str, Any]] = {} + expanded_winners: Dict[str, str] = {} + expanded_stacks: Dict[str, List[Dict[str, Any]]] = {} + + for session_id, session_data in autotune_sessions.items(): + if not session_data: + continue + partitions = _split_session_by_launch_key( + session_id, session_data, launch_by_group_hash + ) + if not partitions: + continue + occurrence_to_group = { + record["occurrence_id"]: record.get("launch_group_hash", "") + for record in session_data.get("launch_occurrences", []) or [] + } + # Deduped coarse compilations define the reference order; shared + # compilations legitimately appear in several partitions. + coarse_compilations = _dedup_compilations_by_hash( + session_data.get("compilations", []) + ) + results = list(session_data.get("autotune_results", []) or []) + legacy_single = session_data.get("autotune_result") + if legacy_single and not results: + results = [legacy_single] + single = len(partitions) == 1 + used_result_indices: Set[int] = set() + # Compilations no partition's launches reference cannot be + # attributed to a key round; promote them to every partition so + # a split never drops them (single-partition sessions already + # keep all coarse compilations). + unattributed_hashes = { + comp.get("payload", {}).get("metadata", {}).get("hash") + for comp in coarse_compilations + } - _globally_referenced_hashes(partitions, launch_by_group_hash) + unattributed_hashes.discard(None) + + for part in partitions: + sub_id = part["sub_session_id"] + sub_data, selected_group = _build_sub_session( + part, + coarse_compilations, + occurrence_to_group, + results, + used_result_indices, + unattributed_hashes, + single, + launch_by_group_hash, + ) + expanded_sessions[sub_id] = sub_data + if selected_group: + expanded_winners[sub_id] = selected_group + expanded_stacks[sub_id] = session_stacks.get(session_id, []) + + return expanded_sessions, expanded_winners, expanded_stacks + + def _generate_autotune_analysis_events( autotune_sessions: Dict[str, Dict[str, Any]], - autotune_winners: Dict[str, str], compilations_by_hash: Dict[str, Any], session_stacks: Dict[str, List[Dict[str, Any]]], launch_by_group_hash: Dict[str, Dict[str, Any]], @@ -118,11 +827,15 @@ def _generate_autotune_analysis_events( """ Generates autotune_analysis events from grouped compilation sessions. + Call-site sessions are first expanded into per-autotune-key + sub-sessions (see _expand_sessions_by_launch_key); every pass below + then operates on sub-sessions. Unsplit sessions keep their plain + session id and behave exactly as before. + Args: autotune_sessions: Dict mapping session_id to - {"compilations": [...], "launch_group_hashes": set([...])}. - autotune_winners: Dict mapping session_id to the selected launch_group_hash - (the winning compilation hash is derived from the launch event). + {"compilations": [...], "launch_group_hashes": set([...]), + "launch_occurrences": [...], "autotune_results": [...]}. compilations_by_hash: Dict containing processed kernel data, used to find output files. session_stacks: Dict mapping session_id to the user call stack. @@ -134,6 +847,15 @@ def _generate_autotune_analysis_events( """ output_events: Dict[str, List[str]] = defaultdict(list) + # Expand each call-site session into per-autotune-key sub-sessions. + # Winners are resolved per sub-session from occurrence records, and + # each sub-session carries its own matched autotune result. + expanded_sessions, expanded_winners, expanded_stacks = ( + _expand_sessions_by_launch_key( + autotune_sessions, session_stacks, launch_by_group_hash + ) + ) + # Pre-compute hash-deduped compilations per session. Cross-PID merge # in parse_single_rank intentionally appends every observation; the # analysis below MUST collapse them by hash before counting "configs" @@ -141,7 +863,7 @@ def _generate_autotune_analysis_events( # _dedup_compilations_by_hash for the rationale. deduped_compilations: Dict[str, List[Dict[str, Any]]] = { sid: _dedup_compilations_by_hash(sd.get("compilations", []) if sd else []) - for sid, sd in autotune_sessions.items() + for sid, sd in expanded_sessions.items() } # First pass: Build hash → groups mapping from sessions with benchmarks @@ -149,7 +871,7 @@ def _generate_autotune_analysis_events( # This allows cached sessions to be associated with all possible groups hash_to_groups: Dict[str, List[List[str]]] = defaultdict(list) - for _session_id, session_data in autotune_sessions.items(): + for _session_id, session_data in expanded_sessions.items(): if not session_data: continue compilation_events = deduped_compilations[_session_id] @@ -176,7 +898,7 @@ def _generate_autotune_analysis_events( hash_to_groups[h].append(compilation_hashes) # Second pass: Generate autotune_analysis events - for session_id, session_data in autotune_sessions.items(): + for session_id, session_data in expanded_sessions.items(): if not session_data: continue @@ -498,7 +1220,7 @@ def _is_tensor_value(val: Any) -> bool: # Resolve winner_compilation_hash from selected launch_group_hash winner_compilation_hash: Optional[str] = None - selected_launch_group_hash = autotune_winners.get(session_id) + selected_launch_group_hash = expanded_winners.get(session_id) if ( selected_launch_group_hash and selected_launch_group_hash in launch_by_group_hash @@ -530,9 +1252,9 @@ def _is_tensor_value(val: Any) -> bool: analysis_event: Dict[str, Any] = { "event_type": "autotune_analysis", "session_id": session_id, - "session_stack": session_stacks.get(session_id, []), + "session_stack": expanded_stacks.get(session_id, []), "name": name, - "selected_hash": autotune_winners.get(session_id), + "selected_hash": expanded_winners.get(session_id), "winner_compilation_hash": winner_compilation_hash, "possible_groups": possible_groups, "compilation_analysis": compilation_analysis, @@ -571,14 +1293,14 @@ def _is_tensor_value(val: Any) -> bool: # 1. Winner run after benchmark (winner_occurrence_ids is not empty) # 2. Cached winner call (benchmark_occurrence_ids is empty) winner_run_counts: Dict[str, int] = defaultdict(int) - for session_id, session_data in autotune_sessions.items(): + for session_id, session_data in expanded_sessions.items(): if not session_data: continue winner_occurrence_ids = session_data.get("winner_occurrence_ids", []) # Count sessions that have winner runs (either after benchmark or cached) if len(winner_occurrence_ids) > 0: # Calculate winner_compilation_hash the same way as in Second pass - selected_launch_group_hash = autotune_winners.get(session_id) + selected_launch_group_hash = expanded_winners.get(session_id) if ( selected_launch_group_hash and selected_launch_group_hash in launch_by_group_hash diff --git a/tritonparse/parse/trace_processor.py b/tritonparse/parse/trace_processor.py index f3d025a..7e84f50 100644 --- a/tritonparse/parse/trace_processor.py +++ b/tritonparse/parse/trace_processor.py @@ -819,9 +819,16 @@ def parse_single_rank( "launch_group_hashes": set(), "benchmark_occurrence_ids": [], "winner_occurrence_ids": [], + # Per-launch records so analysis can split one call-site session + # into per-autotune-key sub-sessions. Each entry maps an + # occurrence back to its launch group. + "launch_occurrences": [], + # One entry per AutotuneListener callback. A list (not a single + # dict) because repeated invocations at the same call site with + # different keys each emit their own result. + "autotune_results": [], } ) - autotune_winners: Dict[str, str] = {} session_stacks: Dict[str, Any] = {} launch_by_group_hash: Dict[str, Dict[str, Any]] = {} next_occurrence_id: int = 0 @@ -981,13 +988,18 @@ def parse_single_rank( autotune_sessions[session_id][ "winner_occurrence_ids" ].append(occurrence_id) + autotune_sessions[session_id]["launch_occurrences"].append( + { + "occurrence_id": occurrence_id, + "launch_group_hash": launch_group_hash, + "is_benchmark": is_benchmark, + } + ) if kernel_hash: kernels_by_hash[kernel_hash]["launches"].append( (parsed_json, global_launch_index) ) - if not is_benchmark and session_id: - autotune_winners[session_id] = launch_group_hash global_launch_index += 1 elif event_type == "autotune": @@ -999,14 +1011,17 @@ def parse_single_rank( ): continue if session_id: - autotune_sessions[session_id]["autotune_result"] = { - "best_config": parsed_json.get("best_config"), - "configs_timings": parsed_json.get("configs_timings"), - "duration": parsed_json.get("duration"), - "cache_hit": parsed_json.get("cache_hit"), - "cache_key": parsed_json.get("cache_key"), - "kernel_name": parsed_json.get("kernel_name"), - } + autotune_sessions[session_id]["autotune_results"].append( + { + "best_config": parsed_json.get("best_config"), + "configs_timings": parsed_json.get("configs_timings"), + "duration": parsed_json.get("duration"), + "cache_hit": parsed_json.get("cache_hit"), + "cache_key": parsed_json.get("cache_key"), + "kernel_name": parsed_json.get("kernel_name"), + "autotune_key": parsed_json.get("autotune_key"), + } + ) if user_stack and session_id not in session_stacks: session_stacks[session_id] = user_stack @@ -1093,7 +1108,6 @@ def parse_single_rank( autotune_events_by_file = _generate_autotune_analysis_events( autotune_sessions, - autotune_winners, kernels_by_hash, session_stacks, launch_by_group_hash, diff --git a/website/public/dedicated_log_triton_trace_findhao__mapped.ndjson.gz b/website/public/dedicated_log_triton_trace_findhao__mapped.ndjson.gz index 81514a4..833dc93 100644 Binary files a/website/public/dedicated_log_triton_trace_findhao__mapped.ndjson.gz and b/website/public/dedicated_log_triton_trace_findhao__mapped.ndjson.gz differ diff --git a/website/public/dedicated_log_triton_trace_inductor__mapped.ndjson.gz b/website/public/dedicated_log_triton_trace_inductor__mapped.ndjson.gz index 547a175..4dfb43e 100644 Binary files a/website/public/dedicated_log_triton_trace_inductor__mapped.ndjson.gz and b/website/public/dedicated_log_triton_trace_inductor__mapped.ndjson.gz differ