TT-Metal

https://docs.tenstorrent.com/tt-metal/latest/
Admin
TT-Metal is Tenstorrent's low-level programming framework for accelerated computing, providing...
Tokens:507,274
Snippets:5,760
Trust Score:10
Update:4 months ago
Show doc for...
Tokens:
Raw
### Run a Programming Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/installing.md.txt

Executes a basic tt-metal programming example to verify the installation.

```sh
python3 -m ttnn.examples.usage.run_op_on_device
```

--------------------------------

### Runtime Arguments Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/eltwise_sfpu.rst.txt

Sets up the runtime arguments for the compute kernel, reader kernel, and writer kernel. The reader and writer get the source and destination addresses and sizes, while the compute kernel knows the expected data sizes.

```cpp
SetRuntimeArgs(program, eltwise_sfpu_kernel_id, core, {n_tiles});
SetRuntimeArgs(program, unary_reader_kernel_id, core, {src0_dram_buffer->address(), n_tiles});
SetRuntimeArgs(program, unary_writer_kernel_id, core, {dst_dram_buffer->address(), n_tiles});
```

--------------------------------

### Program Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/dram_loopback.rst.txt

Sets up the command queue and creates a program object.

```cpp
distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();
Program program = CreateProgram();
```

--------------------------------

### Program Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/custom_sfpi_smoothstep.rst.txt

Host-side setup for the custom SFPI smoothstep example, including device initialization, buffer creation, and kernel setup.

```cpp
// Create a 1x1 mesh on device 0 (same API scales to multi-device meshes)
    constexpr int device_id = 0;
    auto mesh_device = distributed::MeshDevice::create_unit_mesh(device_id);

    // Submit work via the mesh command queue: uploads/downloads and program execution
    distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();
    Program program = CreateProgram();

    // A MeshWorkload is a collection of programs that will be executed on the mesh
    distributed::MeshWorkload workload;
    distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());

    // Configure mesh buffers with single-tile page size
    distributed::DeviceLocalBufferConfig dram_config{
        .page_size = tile_size_bytes,
        .buffer_type = BufferType::DRAM};
    distributed::ReplicatedBufferConfig dram_buffer_config{
        .size = dram_buffer_size};

    // DRAM buffers: single input + one output for unary operation
    auto src0_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());
    auto dst_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());

    // Circular buffers for kernel communication
    CreateCircularBuffer(program, core, /* cb_in0 config */);
    CreateCircularBuffer(program, core, /* cb_out0 config */);

    // Kernels: reader, writer, and custom SFPU compute
    auto reader = CreateKernel(program, "..../read_tiles.cpp", core, DataMovementConfig{...});
    auto writer = CreateKernel(program, "..../write_tile.cpp", core, DataMovementConfig{...});
    auto compute = CreateKernel(program, "..../tiles_smoothstep.cpp", core, ComputeConfig{});
```

--------------------------------

### Device Initialization and Program Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_single_core.rst.txt

Initializes the device, command queue, and program, then sets up matrix dimensions and maps operations to a single Tensix core.

```cpp
// Open device (we use device 0, the first available device)
constexpr int device_id = 0;
std::shared_ptr<distributed::MeshDevice> mesh_device = distributed::MeshDevice::create_unit_mesh(device_id);

// Matrix dimensions (must be divisible by tile dimensions)
constexpr uint32_t M = 640;  // Matrix A height
constexpr uint32_t N = 640;  // Matrix B width
constexpr uint32_t K = 640;  // Shared dimension

// Calculate number of tiles in each dimension
uint32_t Mt = M / TILE_HEIGHT;  // Each tile is 32x32
uint32_t Kt = K / TILE_WIDTH;
uint32_t Nt = N / TILE_WIDTH;

distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();
Program program{};
CoreCoord core({0, 0});  // Single core at position {0, 0}
```

--------------------------------

### Build and Run Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/dram_loopback.rst.txt

Instructions on how to build and run the DRAM loopback example.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples
# To run the example
./build/programming_examples/metal_example_loopback
```

--------------------------------

### Program setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/eltwise_binary.rst.txt

Initializes the Metalium mesh device, command queue, program object, defines the core, and sets constants for tile and buffer sizes.

```cpp
constexpr int device_id = 0;
auto mesh_device = distributed::MeshDevice::create_unit_mesh(device_id);

distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();

Program program = CreateProgram();

constexpr CoreCoord core = {0, 0};
constexpr uint32_t n_tiles = 64;
constexpr uint32_t elements_per_tile = tt::constants::TILE_WIDTH * tt::constants::TILE_HEIGHT;
constexpr uint32_t tile_size_bytes = sizeof(bfloat16) * elements_per_tile;
```

--------------------------------

### Program Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/custom_sfpi_add.rst.txt

Host-side setup for the custom SFPI example, including device initialization, buffer creation, circular buffer allocation, and kernel creation.

```cpp
// Standard device and program setup
    constexpr int device_id = 0;
    auto mesh_device = distributed::MeshDevice::create_unit_mesh(device_id);

    // Submit work via the mesh command queue: uploads/downloads and program execution
    distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();
    // Allocate mesh buffers: two inputs + one output (replicated across mesh)
    auto src0_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());
    auto src1_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());
    auto dst_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());

    Program program = CreateProgram();

    // Create mesh workload for program execution across the mesh
    distributed::MeshWorkload workload;
    distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());

    // Circular buffers for kernel communication
    CreateCircularBuffer(program, core, /* cb_in0 config */);
    CreateCircularBuffer(program, core, /* cb_in1 config */);
    CreateCircularBuffer(program, core, /* cb_out0 config */);

    // Kernels: reader, writer, and custom SFPU compute
    auto reader = CreateKernel(program, "..../read_tiles.cpp", core, DataMovementConfig{...});
    auto writer = CreateKernel(program, "..../write_tile.cpp", core, DataMovementConfig{...});
    auto compute = CreateKernel(program, "..../tiles_add.cpp", core, ComputeConfig{});
```

--------------------------------

### Example Test Class Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/_sources/ttnn/demos.rst.txt

Illustrates how to set up a test class for organizing device and end-to-end performance tests, referencing ResNet50TestInfra.

```python
For the purpose of organizing your tests to be managed for collecting both device and end-to-end performance, please take a look at ``test_ttnn_functional_resnet50.py`` and how the ``ResNet50TestInfra`` class was setup.
```

--------------------------------

### TT-Installer Script

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/installing.md.txt

Download and run the TT-Installer installation script for a quick setup.

```bash
curl -fsSL https://github.com/tenstorrent/tt-installer/releases/latest/download/install.sh -O
chmod +x install.sh
./install.sh --no-install-podman --no-install-metalium-container
```

--------------------------------

### Build and Run Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_single_core.rst.txt

Instructions on how to build and run the matmul_single_core programming example.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples
./build/programming_examples/metal_example_matmul_single_core
```

--------------------------------

### Device Initialization & Program Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_multi_core.rst.txt

C++ code snippet showing device initialization, matrix dimension setup, data generation, and buffer creation for a multi-core matmul operation using the mesh API.

```cpp
// Open mesh device (use device 0)
    constexpr int device_id = 0;
    auto mesh_device = distributed::MeshDevice::create_unit_mesh(device_id);

    // Matrix dimensions (must be divisible by tile dimensions)
    constexpr uint32_t M = 640;  // Matrix A height
    constexpr uint32_t N = 640;  // Matrix B width
    constexpr uint32_t K = 640;  // Shared dimension

    // Calculate number of tiles in each dimension
    uint32_t Mt = M / TILE_HEIGHT;  // Each tile is 32x32
    uint32_t Kt = K / TILE_WIDTH;
    uint32_t Nt = N / TILE_WIDTH;

    std::mt19937 rng(std::random_device{}());
    std::uniform_real_distribution<float> dist(-0.5f, 0.5f);

    std::vector<bfloat16> src0_vec(M * K, 0);  // Matrix A (MxK)
    std::vector<bfloat16> src1_vec(K * N, 0);  // Matrix B (KxN)
    // Fill with random bfloat16 values
    for (bfloat16& v : src0_vec) {
        v = bfloat16(dist(rng));
    }
    for (bfloat16& v : src1_vec) {
        v = bfloat16(dist(rng));
    }

    std::vector<bfloat16> golden_vec(M * N, 0);
    golden_matmul(src0_vec, src1_vec, golden_vec, M, N, K);

    // Convert source matrices to tiled format for device execution
    src0_vec = tilize_nfaces(src0_vec, M, K);
    src1_vec = tilize_nfaces(src1_vec, K, N);

    // Set up mesh command queue, workload, and program
    distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();
    distributed::MeshWorkload workload;
    distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());
    Program program{};

    // Create DRAM buffers for input and output matrices (replicated per device across the mesh)
    constexpr uint32_t single_tile_size = sizeof(bfloat16) * TILE_HEIGHT * TILE_WIDTH;
    distributed::DeviceLocalBufferConfig dram_config{
        .page_size = single_tile_size, .buffer_type = tt_metal::BufferType::DRAM};
    distributed::ReplicatedBufferConfig buffer_config_A{.size = single_tile_size * Mt * Kt};
    distributed::ReplicatedBufferConfig buffer_config_B{.size = single_tile_size * Nt * Kt};
    distributed::ReplicatedBufferConfig buffer_config_C{.size = single_tile_size * Mt * Nt};

    auto src0_dram_buffer = distributed::MeshBuffer::create(buffer_config_A, dram_config, mesh_device.get());
    auto src1_dram_buffer = distributed::MeshBuffer::create(buffer_config_B, dram_config, mesh_device.get());

```

--------------------------------

### Host-side Kernel Execution Setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_single_core.rst.txt

This snippet shows how runtime arguments are set for the reader and writer kernels on the host side, and how input data is uploaded to the device.

```cpp
// Set runtime arguments for kernels
    uint32_t src0_addr = src0_dram_buffer->address();
    uint32_t src1_addr = src1_dram_buffer->address();
    uint32_t dst_addr = dst_dram_buffer->address();

    tt_metal::SetRuntimeArgs(program, reader_id, core, {src0_addr, src1_addr, Mt, Kt, Nt});
    tt_metal::SetRuntimeArgs(program, writer_id, core, {dst_addr, Mt, Nt}); // Note: Writer kernel uses Mt, Nt for output C
    // Don't need to set runtime args for compute kernel, as everything is passed as compile-time args

    // Upload input data to device
    distributed::EnqueueWriteMeshBuffer(cq, src0_dram_buffer, a.data(), false);

```

--------------------------------

### Build and run commands

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/eltwise_binary.rst.txt

Commands to build the programming examples and run the specific eltwise binary example.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples
# To run the example
./build/programming_examples/metal_example_eltwise_binary
```

--------------------------------

### Build Script Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/eltwise_sfpu.rst.txt

Instructions on how to build the programming examples, including the eltwise_sfpu example, by setting environment variables and running build scripts.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh
./build/programming_examples/metal_example_eltwise_sfpu
```

--------------------------------

### Build and Run Profiler Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tools/device_program_profiler.rst.txt

Commands to build the programming examples and run the 'test_full_buffer' example with device profiling enabled.

```bash
cd $TT_METAL_HOME
./build_metal.sh --build-programming-examples
TT_METAL_DEVICE_PROFILER=1 ./build/programming_examples/profiler/test_full_buffer
```

--------------------------------

### GetDefaultDevice Example

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.GetDefaultDevice.html

Example of how to get the default device.

```python
>>> device = ttnn.GetDefaultDevice()

```

--------------------------------

### Example

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.signbit.html

Example of using ttnn.signbit to get the sign bit of each element in a tensor.

```python
# Create a tensor with random values
tensor = ttnn.rand([2, 2], dtype=ttnn.bfloat16, layout=ttnn.TILE_LAYOUT, device=device)

# Test the sign bit of each element
output = ttnn.signbit(tensor)
logger.info(f"Sign bit: {output}")

```

--------------------------------

### Run the example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/examples/custom_sfpi_smoothstep.html

Run the example.

```bash
./build/programming_examples/metal_example_custom_sfpi_smoothstep

```

--------------------------------

### Example

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.conj.html

Example of how to use ttnn.conj to get the complex conjugate of a complex tensor.

```python
# Create a complex tensor
real_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
imag_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
tensor = ttnn.complex_tensor(
    ttnn.Tensor(real_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
    ttnn.Tensor(imag_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
)

# Get the complex conjugate
output = ttnn.conj(tensor, memory_config=ttnn.DRAM_MEMORY_CONFIG)
logger.info(f"Complex conjugate: {output}")

```

--------------------------------

### Build the example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/examples/custom_sfpi_smoothstep.html

Build the example using the `--build-programming-examples` flag.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples

```

--------------------------------

### Get the real part

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.real.html

Example of how to get the real part of a complex tensor using ttnn.real.

```python
# Create a complex tensor
real_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
imag_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
tensor = ttnn.complex_tensor(
    ttnn.Tensor(real_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
    ttnn.Tensor(imag_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
)

# Get the real part
output = ttnn.real(tensor, memory_config=ttnn.DRAM_MEMORY_CONFIG)
logger.info(f"Real part: {output}")

```

--------------------------------

### Get the imaginary part

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.imag.html

Example of how to get the imaginary part of a complex tensor using ttnn.imag.

```python
# Create a complex tensor
real_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
imag_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
tensor = ttnn.complex_tensor(
    ttnn.Tensor(real_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
    ttnn.Tensor(imag_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
)

# Get the imaginary part
output = ttnn.imag(tensor, memory_config=ttnn.DRAM_MEMORY_CONFIG)
logger.info(f"Imaginary part: {output}")

```

--------------------------------

### Build and Run Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/custom_sfpi_add.rst.txt

Instructions on how to build and run the custom SFPI add example.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples
# To run the example
./build/programming_examples/metal_example_custom_sfpi_add
```

--------------------------------

### Start Tracy GUI on macOS

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tools/tracy_profiler.rst.txt

Command to start the Tracy GUI after installation on macOS.

```bash
tracy
```

--------------------------------

### ttnn.topk example

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.topk.html

Example of using ttnn.topk to get the top 32 values and their indices along the last dimension.

```python
# Create tensor
tensor_input = ttnn.rand([1, 1, 32, 64], device=device)

# Apply ttnn.topk() to get top 3 values along dim=1
values, indices = ttnn.topk(tensor_input, k=32, dim=-1, largest=True, sorted=True)
logger.info(f"Topk values: {values}")
logger.info(f"Topk indices: {indices}")

```

--------------------------------

### Install and Launch TT-NN Visualizer

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/_sources/ttnn/tutorials/tutorials/ttnn_visualizer.md.txt

Commands to install and launch the TT-NN Visualizer.

```bash
pip install ttnn-visualizer
ttnn-visualizer
```

--------------------------------

### Scale Mask Softmax In Place Example

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.scale_mask_softmax_in_place.html

This example demonstrates the usage of ttnn.scale_mask_softmax_in_place with a basic setup of input tensor and mask.

```python
# Setup input tensor and mask
input_shape = (1, 1, 32, 32)

attention_mask_t = ttnn.rand(input_shape, dtype=ttnn.bfloat8_b, layout=ttnn.TILE_LAYOUT, device=device)
input_tensor = ttnn.rand(input_shape, dtype=ttnn.bfloat8_b, layout=ttnn.TILE_LAYOUT, device=device)

# Apply in-place scale mask softmax
tt_output = ttnn.scale_mask_softmax_in_place(
    input_tensor=input_tensor,
    scale=1.0,
    mask=attention_mask_t,
)
logger.info(f"Scale Mask Softmax In Place result: {tt_output}")

```

--------------------------------

### Example

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.angle.html

Get the angle (phase) of the complex tensor

```python
# Create a complex tensor
real_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
imag_part = torch.rand([1, 1, 32, 32], dtype=torch.bfloat16)
tensor = ttnn.complex_tensor(
    ttnn.Tensor(real_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
    ttnn.Tensor(imag_part, ttnn.bfloat16).to(ttnn.TILE_LAYOUT).to(device),
)

# Get the angle (phase) of the complex tensor
output = ttnn.angle(tensor, memory_config=ttnn.DRAM_MEMORY_CONFIG)
logger.info(f"Angle: {output}")

```

--------------------------------

### Build and Run Example Program

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/labs/matmul/lab1/lab1.rst.txt

Steps to build and run the example program, demonstrating JIT compilation errors.

```bash
./build_metal.sh
./build/ttnn/examples/example_lab_eltwise_binary
```

--------------------------------

### Slice a tensor

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/api/ttnn.slice.html

Example of slicing a tensor with specified start, end, and step values.

```python
# Create a tensor to slice
input_tensor = ttnn.rand((1, 1, 64, 32), dtype=ttnn.bfloat16, layout=ttnn.Layout.TILE, device=device)

# Slice the tensor
sliced_tensor = ttnn.slice(input_tensor, [0, 0, 0, 0], [1, 1, 64, 16], [1, 1, 2, 1])
logger.info("Sliced Tensor Shape:", sliced_tensor.shape)  # Sliced Tensor Shape: Shape([1, 1, 32, 16])
```

--------------------------------

### Build and Run Example

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/examples/eltwise_sfpu.html

Instructions on how to build and run the eltwise_sfpu example.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh
./build/programming_examples/metal_example_eltwise_sfpu

```

--------------------------------

### TT-Installer with Specific Versions

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/installing.md.txt

Example of using flags to specify dependency versions with the TT-Installer script.

```bash
./install.sh \
  --smi-version=v3.0.38 \
  --fw-version=19.2.0 \
  --kmd-version=2.5.0 \
  --no-install-podman \
  --no-install-metalium-container
```

--------------------------------

### Set Up Environment for Models

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/installing.md.txt

Execute commands to set up the environment for using pre-built models.

```sh
export PYTHONPATH=$(pwd)
pip install -r tt_metal/python_env/requirements-dev.txt
sudo apt-get install cpufrequtils
sudo cpupower frequency-set -g performance
```

--------------------------------

### Launching the Program

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/dram_loopback.rst.txt

Code snippet demonstrating how to create a MeshWorkload and enqueue it for execution on the mesh.

```cpp
distributed::MeshWorkload workload;
    distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());
    workload.add_program(device_range, std::move(program));
    distributed::EnqueueMeshWorkload(cq, workload, /*blocking=*/false);
    distributed::Finish(cq);
    // Equivalently, we could have done:
    // distributed::EnqueueMeshWorkload(cq, workload, /*blocking=*/true);
```

--------------------------------

### Build and Run Instructions

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_multi_core.rst.txt

Instructions on how to build the programming examples and run the multi-core matmul executable.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples
./build/programming_examples/metal_example_matmul_multi_core
```

--------------------------------

### Running Tutorials

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/tutorials.html

Example of how to run tutorials using Python 3 and a standalone script.

```bash
$ python3 --version
Python 3.10.12
$ python3 example.py
...

```

--------------------------------

### Reader Kernel for Multi-Core Matmul

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_multi_core.rst.txt

The reader kernel is responsible for reading input data from DRAM buffers and pushing it into circular buffers. It needs to know the number of tiles to read and the starting tile index for each core. It calculates the correct offset in the DRAM buffer based on the starting tile index and the dimensions of the input matrices.

```cpp
void kernel_main() {
    uint32_t src0_addr = get_arg_val<uint32_t>(0);
    uint32_t src1_addr = get_arg_val<uint32_t>(1);
    uint32_t Mt = get_arg_val<uint32_t>(2);
    uint32_t Kt = get_arg_val<uint32_t>(3);
    uint32_t Nt = get_arg_val<uint32_t>(4);
    uint32_t output_tile_start_id = get_arg_val<uint32_t>(5); // Starting tile ID for this core
    uint32_t num_output_tiles = get_arg_val<uint32_t>(6); // Number of output tiles to read

    constexpr uint32_t cb_id_in0 = tt::CBIndex::c_0;
    constexpr uint32_t cb_id_in1 = tt::CBIndex::c_1;

    const uint32_t in0_tile_bytes = get_tile_size(cb_id_in0);
    const uint32_t in1_tile_bytes = get_tile_size(cb_id_in1);

    constexpr auto a_args = TensorAccessorArgs<0>();
    const auto a = TensorAccessor(a_args, src0_addr, in0_tile_bytes);

    constexpr auto b_args = TensorAccessorArgs<a_args.next_compile_time_args_offset()>();
    const auto b = TensorAccessor(b_args, src1_addr, in1_tile_bytes);

    // Loop through the output tiles assigned to this core
    for (uint32_t output_tile = 0; output_tile < num_output_tiles; ++output_tile) {
        uint32_t current_tile_id = output_tile_start_id + output_tile;

        // Calculate the output tile position in the grid
        uint32_t out_row = current_tile_id / Nt;
        uint32_t out_col = current_tile_id % Nt;

        // Read all K tiles for this output position. Same inner loop as in the single core example.
        for (uint32_t k = 0; k < Kt; k++) {
            uint32_t tile_A = out_row * Kt + k;
            {
                cb_reserve_back(cb_id_in0, 1);
                uint32_t l1_write_addr_in0     = get_write_ptr(cb_id_in0);
                noc_async_read_tile(tile_A, a, l1_write_addr_in0);
                noc_async_read_barrier();
                cb_push_back(cb_id_in0, 1);
            }

            uint32_t tile_B = k * Nt + out_col;
            {
                cb_reserve_back(cb_id_in1, 1);
                uint32_t l1_write_addr_in1 = get_write_ptr(cb_id_in1);
                noc_async_read_tile(tile_B, b, l1_write_addr_in1);
                noc_async_read_barrier();
                cb_push_back(cb_id_in1, 1);
            }
        }
    }
}
```

--------------------------------

### Writer Kernel for Partitioned Work

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/tt_metal/examples/matmul_multi_core.rst.txt

This C++ kernel code demonstrates how to partition work for a writer kernel in Metalium, where each core processes a specific number of tiles and writes them to the DRAM buffer starting from a given index.

```cpp
    void kernel_main() {
        uint32_t dst_addr = get_arg_val<uint32_t>(0);
        uint32_t num_tiles = get_arg_val<uint32_t>(1); // Number of tiles to write
        uint32_t start_id = get_arg_val<uint32_t>(2);  // Starting tile ID for this core

        constexpr uint32_t cb_id_out = tt::CBIndex::c_16;

        const uint32_t tile_bytes = get_tile_size(cb_id_out);

        constexpr auto c_args = TensorAccessorArgs<0>();
        const auto c = TensorAccessor(c_args, dst_addr, tile_bytes);

        // Each core writes only its assigned tiles
        for (uint32_t i = 0; i < num_tiles; ++i) {
            cb_wait_front(cb_id_out, 1);
            uint32_t l1_read_addr = get_read_ptr(cb_id_out);
            // Write to the correct offset based on start_id
            noc_async_write_tile(i + start_id, c, l1_read_addr);

```

--------------------------------

### Build and Run Commands

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/examples/eltwise_binary.html

Commands to build the programming examples and run the specific eltwise binary example.

```bash
export TT_METAL_HOME=</path/to/tt-metal>
./build_metal.sh --build-programming-examples
# To run the example
./build/programming_examples/metal_example_eltwise_binary

```

--------------------------------

### Program setup

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/examples/custom_sfpi_smoothstep.html

The host-side setup for this custom SFPI example follows the standard pattern: mesh device initialization, DRAM buffer creation, circular buffer allocation for kernel communication, and kernel creation. Unlike binary operations that require two input buffers, smoothstep is a unary operation requiring only a single input buffer (`src0_dram_buffer`) plus the output buffer (`dst_dram_buffer`). Correspondingly, we need only one input circular buffer (`cb_in0`) and one output buffer (`cb_out0`).

```cpp
// Create a 1x1 mesh on device 0 (same API scales to multi-device meshes)
constexpr int device_id = 0;
auto mesh_device = distributed::MeshDevice::create_unit_mesh(device_id);

// Submit work via the mesh command queue: uploads/downloads and program execution
distributed::MeshCommandQueue& cq = mesh_device->mesh_command_queue();
Program program = CreateProgram();

// A MeshWorkload is a collection of programs that will be executed on the mesh
distributed::MeshWorkload workload;
distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());

// Configure mesh buffers with single-tile page size
distributed::DeviceLocalBufferConfig dram_config{
    .page_size = tile_size_bytes,
    .buffer_type = BufferType::DRAM};
distributed::ReplicatedBufferConfig dram_buffer_config{
    .size = dram_buffer_size};

// DRAM buffers: single input + one output for unary operation
auto src0_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());
auto dst_dram_buffer = distributed::MeshBuffer::create(dram_buffer_config, dram_config, mesh_device.get());

// Circular buffers for kernel communication
CreateCircularBuffer(program, core, /* cb_in0 config */);
CreateCircularBuffer(program, core, /* cb_out0 config */);

// Kernels: reader, writer, and custom SFPU compute
auto reader = CreateKernel(program, "..../read_tiles.cpp", core, DataMovementConfig{...});
auto writer = CreateKernel(program, "..../write_tile.cpp", core, DataMovementConfig{...});
auto compute = CreateKernel(program, "..../tiles_smoothstep.cpp", core, ComputeConfig{});

```

--------------------------------

### Full Example and Output

Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/ttnn/tutorials/tutorials/ttnn_add_tensors.html

Combines all previous examples into a complete script that opens a device, creates tensors, performs addition, and logs the output. The expected output is a tensor filled with 3s.

```shell
$ python3 $TT_METAL_HOME/ttnn/tutorials/basic_python/ttnn_add_tensors.py
```

--------------------------------

### Install Latest Wheel

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/_sources/installing.md.txt

Install the ttnn wheel using pip.

```sh
pip install ttnn
```
Tokens:
Raw