Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

GraphCapture doesn't recognise CBs with globally allocated addresses #12624

Closed
derdeljanTT opened this issue Sep 12, 2024 · 1 comment
Closed
Assignees
Labels
bug Something isn't working forge

Comments

@derdeljanTT
Copy link
Contributor

Describe the bug
Graph capture doesn't recognise that CBs with globally allocated addresses will share space with the corresponding sharded input / output operand and no additional bytes will be allocated.

To Reproduce
Go to branch derdeljan/graph-capture-cbs-repro, build tt-metal and run ./build_Debug/test/ttnn/unit_tests_ttnn --gtest_filter=MlirInterfaceTests/EltwiseUnaryOpInterfaceTestFixture.MlirInterfaceTest/0 (on a lab machine). Repro test definition is in the file tests/ttnn/unit_tests/gtests/test_graph_capture_repro.cpp.

We can see from the logs:

Always | INFO     | CB In0 has the same address as the sharded input 0
Always | INFO     | CB Out0 has the same address as the sharded output 0

that the two CBs being allocated in unary sharded program factory have the sam address as the corresponding operand, however graph capture reports as if the new space will be allocated:

circular buffers:
393216
393216

Full trace:

                 Always | INFO     | Trace: [
    {
        "connections": [
            1,
            19
        ],
        "counter": 0,
        "node_type": "capture_start",
        "params": {}
    },
    {
        "connections": [
            3,
            4,
            17
        ],
        "counter": 1,
        "node_type": "function_start",
        "params": {
            "inputs": "1",
            "name": "ttnn::relu"
        }
    },
    {
        "connections": [
            1
        ],
        "counter": 2,
        "node_type": "tensor",
        "params": {
            "shape": "ttnn.Shape([3, 1, 1024, 1024])",
            "tensor_id": "0"
        }
    },
    {
        "connections": [
            2,
            5
        ],
        "counter": 3,
        "node_type": "buffer",
        "params": {
            "layout": "HEIGHT_SHARDED",
            "size": "6291456",
            "type": "L1"
        }
    },
    {
        "connections": [
            6,
            16
        ],
        "counter": 4,
        "node_type": "function_start",
        "params": {
            "inputs": "8",
            "name": "ttnn::prim::unary"
        }
    },
    {
        "connections": [
            4
        ],
        "counter": 5,
        "node_type": "tensor",
        "params": {
            "shape": "ttnn.Shape([3, 1, 1024, 1024])",
            "tensor_id": "1"
        }
    },
    {
        "connections": [
            7,
            12,
            13,
            14,
            15
        ],
        "counter": 6,
        "node_type": "function_start",
        "params": {
            "inputs": "2",
            "name": "UnaryDeviceOperation"
        }
    },
    {
        "connections": [
            8,
            9,
            10
        ],
        "counter": 7,
        "node_type": "function_start",
        "params": {
            "inputs": "5",
            "name": "tt::tt_metal::create_device_tensor"
        }
    },
    {
        "connections": [
            11,
            11,
            11,
            18
        ],
        "counter": 8,
        "node_type": "buffer",
        "params": {
            "layout": "HEIGHT_SHARDED",
            "size": "6291456",
            "type": "L1"
        }
    },
    {
        "connections": [
            8
        ],
        "counter": 9,
        "node_type": "buffer_allocate",
        "params": {
            "address": "1866087022",
            "layout": "HEIGHT_SHARDED",
            "size": "6291456",
            "type": "L1"
        }
    },
    {
        "connections": [
            11
        ],
        "counter": 10,
        "node_type": "function_end",
        "params": {
            "name": "tt::tt_metal::create_device_tensor"
        }
    },
    {
        "connections": [],
        "counter": 11,
        "node_type": "tensor",
        "params": {
            "shape": "ttnn.Shape([3, 1, 1024, 1024])",
            "tensor_id": "2"
        }
    },
    {
        "connections": [
            6
        ],
        "counter": 12,
        "node_type": "circular_buffer_deallocate_all",
        "params": {}
    },
    {
        "connections": [],
        "counter": 13,
        "node_type": "circular_buffer_allocate",
        "params": {
            "address": "0",
            "core_range_set": "{[(x=0,y=0) - (x=3,y=3)]}",
            "size": "393216"
        }
    },
    {
        "connections": [],
        "counter": 14,
        "node_type": "circular_buffer_allocate",
        "params": {
            "address": "0",
            "core_range_set": "{[(x=0,y=0) - (x=3,y=3)]}",
            "size": "393216"
        }
    },
    {
        "connections": [
            11
        ],
        "counter": 15,
        "node_type": "function_end",
        "params": {
            "name": "UnaryDeviceOperation"
        }
    },
    {
        "connections": [
            11
        ],
        "counter": 16,
        "node_type": "function_end",
        "params": {
            "name": "ttnn::prim::unary"
        }
    },
    {
        "connections": [
            18,
            20
        ],
        "counter": 17,
        "node_type": "function_end",
        "params": {
            "name": "ttnn::relu"
        }
    },
    {
        "connections": [],
        "counter": 18,
        "node_type": "tensor",
        "params": {
            "shape": "ttnn.Shape([3, 1, 1024, 1024])",
            "tensor_id": "3"
        }
    },
    {
        "connections": [
            8
        ],
        "counter": 19,
        "node_type": "buffer_deallocate",
        "params": {
            "layout": "HEIGHT_SHARDED",
            "size": "0",
            "type": "L1"
        }
    },
    {
        "connections": [],
        "counter": 20,
        "node_type": "capture_end",
        "params": {}
    }
]

Expected behavior
Graph capture should recognise that no new space will be allocated for the above-mentioned CBs.

@derdeljanTT derdeljanTT added the bug Something isn't working label Sep 12, 2024
derdeljanTT added a commit that referenced this issue Nov 1, 2024
derdeljanTT added a commit that referenced this issue Nov 4, 2024
derdeljanTT added a commit that referenced this issue Nov 4, 2024
* #12624: Fix tracking of globally allocated CBs in graph capture

* #12624: Update graph tracing docs
ct-clmsn pushed a commit to ct-clmsn/tt-metal that referenced this issue Nov 12, 2024
…pture (tenstorrent#14521)

* tenstorrent#12624: Fix tracking of globally allocated CBs in graph capture

* tenstorrent#12624: Update graph tracing docs
@mbezuljTT
Copy link
Contributor

#14521 merged, closing this issue

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working forge
Projects
None yet
Development

No branches or pull requests

4 participants