TT-Metalium

https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/
Admin
TT-Metalium is a programming framework for Tenstorrent AI accelerators, providing APIs and tools to...
Tokens:246,249
Snippets:1,619
Trust Score:10
Update:10 months ago
Show doc for...
Tokens:
Raw
### Program Setup for Custom SFPI Smoothstep

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

Host-side setup for the custom SFPI smoothstep example, including mesh device initialization, buffer creation, and kernel compilation. This snippet demonstrates buffer configuration and kernel instantiation for a unary operation.

```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{});
```

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

### Install TT-Installer Script

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

Installs Tenstorrent software stack using the TT-Installer script. This is the recommended method for a quick setup. It downloads, makes executable, and runs the installation script. Flags can be used to specify dependency versions and skip certain installations like Podman or Metalium containers.

```sh
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
```

```sh
./install.sh \
  --smi-version=v3.0.17 \
  --fw-version=18.3.0 \
  --kmd-version=1.34 \
  --no-install-podman \
  --no-install-metalium-container
```

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

### Build and Run tt-metal Programming Examples (Bash)

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

Commands to build the tt-metal project and execute a specific programming example. This includes setting the TT_METAL_HOME environment variable and running the compiled 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
```

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

### Setup Circular Buffers for Eltwise SFPU Example

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

This C++ code demonstrates the setup of circular buffers for input and output in the tt-metal SFPU example. It configures two circular buffers with a page size equal to the tile size, allowing for overlapping data movement and compute operations.

```cpp
// Allocate 2 circular buffers for input and output.
    constexpr uint32_t src0_cb_index = tt::CBIndex::c_0;
    constexpr uint32_t num_input_tiles = 2;
    CircularBufferConfig cb_src0_config = 
        CircularBufferConfig(num_input_tiles * tile_size_bytes, {{src0_cb_index, tt::DataFormat::Float16_b}})
            .set_page_size(src0_cb_index, tile_size_bytes);
    CBHandle cb_src0 = tt_metal::CreateCircularBuffer(program, core, cb_src0_config);

    constexpr uint32_t output_cb_index = tt::CBIndex::c_16;
    constexpr uint32_t num_output_tiles = 2;
    CircularBufferConfig cb_output_config = 
        CircularBufferConfig(num_input_tiles * tile_size_bytes, {{output_cb_index, tt::DataFormat::Float16_b}})
            .set_page_size(output_cb_index, tile_size_bytes);
    CBHandle cb_output = tt_metal::CreateCircularBuffer(program, core, cb_output_config);
```

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

### C++: Device Initialization and Program Setup for Matmul

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

This C++ code snippet illustrates the initialization of a Tensix device and the setup of a program for matrix multiplication on a single core. It translates matrix dimensions into tile-based dimensions suitable for hardware operation and defines the target core's coordinates.

```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}
```

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

### Setup TT-Metal Program and Command Queue (C++)

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

Sets up the necessary components for executing programs on the TT-Metal device. This includes obtaining a mesh command queue for managing asynchronous operations and creating a Program object to hold the kernels to be executed.

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

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

### Building and Running tt-metal Profiling Example (Shell)

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

These shell commands demonstrate how to build the programming examples for tt-metal and then run the `test_full_buffer` example with the device profiler enabled. The `TT_METAL_DEVICE_PROFILER=1` environment variable is crucial for activating the profiler.

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

```

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

### Run Custom SFPI Smoothstep Example

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

Executes the compiled custom SFPI smoothstep programming example. This command initiates the smoothstep computation on the Tenstorrent hardware.

```bash
./build/programming_examples/metal_example_custom_sfpi_smoothstep
```

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

### Build and Run DRAM Loopback Example

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

This bash script demonstrates how to build and run the DRAM loopback example. It involves setting the TT_METAL_HOME environment variable, building the programming examples with a specific flag, and then executing the compiled loopback executable.

```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
```

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

### Build and Run Eltwise SFPU Example

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

This bash script demonstrates how to build and execute the Eltwise SFPU programming example for tt-metal. It involves setting the TT_METAL_HOME environment variable, running the build script, and then executing the compiled example.

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

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

### Run Profiler Example

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

Executes the 'test_full_buffer' programming example after building it with profiling enabled. Requires the TT_METAL_DEVICE_PROFILER environment variable to be set.

```bash
TT_METAL_DEVICE_PROFILER=1 ./build/programming_examples/profiler/test_full_buffer
```

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

### C++: tt-metal Host-Side Program Setup for Vector Addition

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

This C++ code snippet illustrates the host-side setup for a custom SFPI vector addition example in tt-metal. It includes device initialization, DRAM buffer creation for inputs and outputs, circular buffer allocation, and kernel creation for reader, writer, and compute operations.

```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{});
```

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

### Initialize tt-metal Mesh Device and Program (C++)

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

Initializes the Metalium mesh device, obtains the command queue, creates a program object, defines the core coordinates, and sets constants for tile and buffer sizes. This setup is essential for subsequent operations.

```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;
```

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

### Bash Script: Build and Run Matmul Single Core Example

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

This bash script demonstrates how to build and execute the matmul_single_core programming example. It requires setting the TT_METAL_HOME environment variable and uses build scripts to compile the example, resulting in an executable for testing matrix multiplication on the device.

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

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

### Build TT-Metal Library

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

Build the TT-Metal library from source. Two options are provided: using a convenient build script (`build_metal.sh`) or a manual build process with CMake for more control over build options. The manual build involves creating a build directory, configuring with CMake, and then using Ninja to build and install the library.

```bash
./build_metal.sh
```

```bash
mkdir build
cd build
cmake .. -G Ninja -DCMAKE_BUILD_TYPE=RelWithDebugInfo -DCMAKE_CXX_COMPILER=<your compiler>
ninja
ninja install
```

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

### Clone TT-Metal Repository

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

Clone the TT-Metal repository from GitHub, including all submodules. This is the first step for installing TT-Metalium from source, which is recommended for developers who want close access to the metal and source code, especially for running demo models.

```bash
git clone https://github.com/tenstorrent/tt-metal.git --recurse-submodules
```

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

### Build and Run Multi-Core Matmul Example

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

Demonstrates how to build and execute the multi-core matmul example. This involves adding a build flag to the build script or CMake command to include programming examples. The compiled executable is then run to perform the multi-core matrix multiplication.

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

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

### Build Programming Examples with Profiler

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

Builds the Metalium project including programming examples, with profiling support enabled. This command is necessary to run examples like 'test_full_buffer'.

```bash
./build_metal.sh --build-programming-examples
```

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

### C++: Host-side Setup for Vector Addition with SFPI

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

This C++ code outlines the host-side setup for performing vector addition using SFPI. It includes device initialization, buffer creation for inputs and outputs, circular buffer allocation for kernel communication, and the creation of reader, writer, and compute kernels. This serves as a template for custom SFPI operations.

```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{});

```

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

### Initialize Mesh Device and Prepare for Matrix Multiplication

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

Initializes a distributed mesh device, defines matrix dimensions, creates input and output buffers in DRAM, and sets up the program for execution. This example assumes TILE_HEIGHT and TILE_WIDTH are defined elsewhere.

```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
// Note: This operation is also shown in a separate snippet below.
// 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());
auto dst_dram_buffer = distributed::MeshBuffer::create(buffer_config_C, dram_config, mesh_device.get());

```

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

### Program Setup for Custom SFPI Smoothstep Kernel

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

Sets up the host-side environment for a custom SFPI kernel, including device initialization, buffer creation, and kernel compilation. It handles mesh device setup, DRAM and circular buffer allocation for a unary operation (smoothstep), and defines reader, writer, and compute kernels.

```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{});

```

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

### Build Custom SFPI Smoothstep Example

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

Builds the custom SFPI smoothstep programming example using the tt-metal build script. Requires setting the TT_METAL_HOME environment variable.

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

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

### Bash Script: Build and Run tt-metal Custom SFPI Add Example

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

This bash script demonstrates how to build and run the custom SFPI vector addition example for tt-metal. It involves setting the TT_METAL_HOME environment variable and executing build and run commands.

```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
```

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

### Install TT-NN Python Package from Binaries

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

Installs the TT-NN Python package using pip. This method provides quick access to TT-NN Python APIs and pre-built AI models. It requires Linux with glibc 2.34 or newer. Post-installation setup for models includes setting PYTHONPATH, installing requirements, and configuring CPU governor for performance.

```sh
pip install ttnn
```

```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
```

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

### Pull and Run TT-Metalium Docker Image

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

Downloads the latest TT-Metalium release image from the Docker registry and runs it in an interactive container. This is a quick way to access APIs and start running AI models. The `--device /dev/tenstorrent` flag is crucial for hardware access within the container.

```sh
docker pull ghcr.io/tenstorrent/tt-metal/tt-metalium-ubuntu-22.04-release-amd64:latest-rc
docker run -it --rm --device /dev/tenstorrent ghcr.io/tenstorrent/tt-metal/tt-metalium-ubuntu-22.04-release-amd64:latest-rc bash
```

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

### Create Data Movement and Compute Kernels

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

Illustrates the creation of reader, writer, and compute kernels for data processing on Tensix cores. It shows how to pass arguments and specify configurations like processor, NoC, and math fidelity.

```cpp
std::vector<uint32_t> reader_args;
TensorAccessorArgs(*src0_dram_buffer->get_backing_buffer()).append_to(reader_args);
TensorAccessorArgs(*src1_dram_buffer->get_backing_buffer()).append_to(reader_args);
auto reader = CreateKernel(
    program,
    "eltwise_binary/kernels/dataflow/read_tiles.cpp",
    core,
    DataMovementConfig{.processor = DataMovementProcessor::RISCV_0, .noc = NOC::RISCV_0_default, .compile_args = reader_args});
std::vector<uint32_t> writer_args;
TensorAccessorArgs(*dst_dram_buffer->get_backing_buffer()).append_to(writer_args);
auto writer = CreateKernel(
    program,
    "eltwise_binary/kernels/dataflow/write_tile.cpp",
    core,
    DataMovementConfig{.processor = DataMovementProcessor::RISCV_1, .noc = NOC::RISCV_1_default, .compile_args = writer_args});
auto compute = CreateKernel(
    program,
    "eltwise_binary/kernels/compute/tiles_add.cpp",
    core,
    ComputeConfig{.math_fidelity = MathFidelity::HiFi4});
```

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

### Reader Kernel Runtime Arguments Setup in C++

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/examples/matmul_multi_core_optimizations/data_reuse

Sets up the runtime arguments for the reader kernel. This vector configures memory addresses, tile start IDs, strides, block dimensions, and other parameters for reading data from DRAM buffers into L1. It includes arguments for both input tensors (in0 and in1) and parameters for block and batch processing.

```cpp
std::vector<uint32_t> mm_reader_args = {
    (std::uint32_t)  src0_dram_buffer->address(), // in0_tensor_addr
    (std::uint32_t)  Kt * per_core_M * output_idx_y, // in0_tensor_start_tile_id
    (std::uint32_t)  1, // in0_tensor_stride_w
    (std::uint32_t)  Kt, // in0_tensor_stride_h
    (std::uint32_t)  in0_block_w, // in0_tensor_next_block_stride

    (std::uint32_t)  in0_block_w, // in0_block_w
    (std::uint32_t)  per_core_M, // in0_block_h
    (std::uint32_t)  in0_block_w * per_core_M, //in0_block_num_tiles

    (std::uint32_t)  src1_dram_buffer->address(), // in1_tensor_addr
    (std::uint32_t)  per_core_N * output_idx_x, //in1_tensor_start_tile_id
    (std::uint32_t)  1, // in1_tensor_stride_w
    (std::uint32_t)  Nt, // in1_tensor_stride_h
    (std::uint32_t)  in0_block_w * Nt, //in1_tensor_next_block_stride

    (std::uint32_t)  per_core_N, // in1_block_w
    (std::uint32_t)  in0_block_w, //in1_block_h
    (std::uint32_t)  per_core_N * in0_block_w, // in1_block_num_tiles

    (std::uint32_t)  Kt / in0_block_w, // num_blocks

    (std::uint32_t)  Mt * Kt, // MtKt
    (std::uint32_t)  Kt * Nt, // KtNt
    (std::uint32_t)  B, // batch
    (std::uint32_t)  bcast_batch // bcast_B
};

```

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

### Initialize rsub_tile Kernel (C++)

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/compute/rsub_tile

Initializes the rsub_tile kernel. Refer to the documentation for any_init for further details.

```c++
void ckernel::rsub_tile_init() {
    // Refer to documentation for any_init.
}
```

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

### ckernel::compute_kernel_hw_startup (3 parameters)

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/compute/compute_kernel_hw_startup

Performs the required hardware initialization for all subsequent operations in the compute kernel. This function should be called exactly once at the very beginning of the kernel, before any operation-specific initialization functions.

```APIDOC
## ckernel::compute_kernel_hw_startup (3 parameters)

### Description
Performs the required hardware initialization for all subsequent operations in the compute kernel. This function should be called exactly once at the very beginning of the kernel, before any operation-specific initialization functions (such as reduce_init, tilize_init, etc.). The circular buffer (CB) IDs provided to this function must match those used in the next operation-specific initialization function. If the operands for the next operation require a different data format than what was configured here, you must call one of the reconfig_data_format functions before proceeding with the next initialization. Similarly, if the next operation requires different properties (such as tile or face dimensions), you must ensure that the same CB IDs are used as in this function.

NOTE: This function performs MMIO writes, which are slow and almost exclusively require the idle state of the execution units that should be configured (PACK, MATH, UNPACK, CFG, etc.). This is why it is unsafe to call this function in the middle of a kernel execution. This function should be called only once at the beginning of the kernel, before any other calls to Compute API are made (either init or other). Calling this function after other API calls may lead cause race conditions and undefined behavior which can be hard to debug.

### Method
N/A (Function Call)

### Parameters
#### Path Parameters
None

#### Query Parameters
None

#### Request Body
None

### Request Example
```c++
ckernel::compute_kernel_hw_startup(icb0, icb1, ocb);
```

### Response
#### Success Response (N/A)
None

#### Response Example
None
```

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

### Host-Side Compute Kernel Configuration (C++)

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/advanced_topics/compute_engines_and_dataflow_within_tensix

This C++ example shows how to configure the Dst register set's behavior on the host when creating a compute kernel. `fp32_dest_acc_en` controls data width (16-bit or 32-bit), and `dst_full_sync_en` toggles double-buffering to optimize performance by overlapping computation and packing.

```cpp
// Example host-side kernel configuration
auto kernel_id = tt::tt_metal::CreateKernel(
    program,
    "path/to/your/compute/kernel.cpp",
    core,
    tt::tt_metal::ComputeConfig{
        .fp32_dest_acc_en = true, // Use 32-bit Dst registers
        .dst_full_sync_en = false  // Enable double-buffering
    }
);

// Number of Dst registers can be checked using ttnn::get_dest_reg_count(const ComputeKernelConfig&)

```

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

### Initialize bitwise_and_tile Kernel

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/compute/bitwise_and_tile

Initializes the bitwise_and_tile kernel. Refer to the documentation for any_init for detailed information.

```c++
void ckernel::bitwise_and_tile_init() {
    // Refer to documentation for any_init.
}
```

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

### Python: Writing Custom SFPU Operations with SFPI

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/index

Guides users on creating custom SFPU operations using the SFPI (Scalar Floating Point Interface) library within TT-Metal. Includes examples for vector addition and smoothstep.

```python
import tt_metal
# Example: Vector Addition using SFPI
def vector_add_sfpi(device):
    program = tt_metal.Program.CreateProgram()
    # ... (program setup, data movement kernels, SFPI compute kernel definition) ...
    tt_metal.Program.ExecuteProgram(program, device)

# Example: Smoothstep using SFPI
def smoothstep_sfpi(device):
    program = tt_metal.Program.CreateProgram()
    # ... (program setup, data movement kernels, custom SFPI smoothstep implementation) ...
    tt_metal.Program.ExecuteProgram(program, device)

```

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

### Configure Circular Buffer in tt-metal (C++)

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

Configures a circular buffer (CB) for inter-kernel communication. This setup specifies the total size, data format, and page size for the circular buffer, using CBIndex `c_0` for source data.

```cpp
    constexpr uint32_t tiles_per_cb = 2;
    tt::CBIndex src0_cb_index = tt::CBIndex::c_0;
    CircularBufferConfig c0_cfg = CircularBufferConfig(
        /*total_size=*/tiles_per_cb * tile_size_bytes,
        /*data_format_spec=*/{{src0_cb_index, tt::DataFormat::Float16_b}})
        .set_page_size(src0_cb_index, tile_size_bytes);
```

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

### Initialize Core Grid and Ranges for Data Distribution

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

Initializes the core grid configuration by defining the starting core coordinates and calculating the core range (width and height) based on the desired number of blocks and cores. This setup is crucial for distributing tile work and assigning broadcast roles.

```cpp
CoreCoord start_core = {0, 0};
CoreCoord core_range = bmm_op_utils::get_core_range(num_blocks_y, num_blocks_x, num_cores_y, num_cores_x);

uint32_t start_core_x = start_core.x;
uint32_t start_core_y = start_core.y;
uint32_t num_cores_c = core_range.x;  // Core count along x-axis
uint32_t num_cores_r = core_range.y;  // Core count along y-axis

```

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

### Writer Kernel for Partitioned Work in tt-metal

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

The writer kernel is responsible for writing output tiles to the DRAM buffer. It accepts parameters for the destination address, the total number of tiles to write, and the starting tile index for the current core. Each core writes only its assigned tiles to the correct offset.

```c++
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);
        noc_async_write_barrier();
        cb_pop_front(cb_id_out, 1);
    }
}

```

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

### Get Circular Buffer Read Pointer (C/C++)

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/circular_buffers/get_read_ptr

Retrieves a pointer to the start of a circular buffer's memory block. This function is valid only after a cb_wait_front call and before a cb_pop_front call. The amount of available data is determined by the number of tiles requested in the preceding cb_wait_front call. The operand specifies the index of the circular buffer.

```c++
uint32_t get_read_ptr(uint32_t operand)
```

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

### Compute Kernel Hardware Startup API

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/compute/compute

API for hardware startup of compute kernels.

```APIDOC
## Compute Kernel Hardware Startup API

### Description
API to initiate hardware startup for compute kernels.

### Method
N/A (Function calls within kernels)

### Endpoint
N/A

### Parameters
N/A

### Request Example
N/A

### Response
N/A

#### Function:
- `compute_kernel_hw_startup`
```

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

### Partition Work in Metalium Writer Kernel (C++)

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

Illustrates a writer kernel in Metalium designed for partitioned work distribution. Each core processes a specific subset of output tiles, defined by the number of tiles and a starting tile ID passed as arguments. This ensures each core contributes only its assigned portion to the final output buffer.

```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);

```

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

### Configure and Execute Kernels on Tenstorrent Device

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

This snippet demonstrates setting runtime arguments for reader and writer kernels, uploading input data to the device, enqueuing a mesh workload for execution, and reading the resulting data back to host memory. It covers essential steps for kernel execution on the Tenstorrent mesh 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);
distributed::EnqueueWriteMeshBuffer(cq, src1_dram_buffer, b.data(), false);

// execute program, and read results
distributed::MeshWorkload workload;
distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());
workload.add_program(device_range, std::move(program));
distributed::EnqueueMeshWorkload(cq, workload, false);
distributed::EnqueueReadMeshBuffer(cq, output.data(), dst_dram_buffer, true);

```

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

### Install Latest TT-Metalium Wheel

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/index

Installs the latest pre-compiled wheel package for TT-Metalium. This is a recommended installation method for users of Tenstorrent models.

```bash
pip install tt_metal-latest-wheel
```

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

### lez_tile_init API

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/compute/lez_tile

Initializes the lez_tile functionality. Refer to any_init documentation for details.

```APIDOC
## ckernel::lez_tile_init()

### Description
Initializes the lez_tile functionality. Please refer to documentation for any_init.

### Method
N/A (CKernel function)

### Endpoint
N/A (CKernel function)

### Parameters
None

### Request Example
```json
{}
```

### Response
#### Success Response (N/A)
N/A

#### Response Example
```json
{}
```
```

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

### Reader Kernel for Partitioned Work in tt-metal

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

The reader kernel fetches input data from DRAM buffers and loads it into circular buffers. It requires input matrix dimensions (Mt, Kt, Nt), the starting tile index for the core, and the number of output tiles to read. It calculates the correct offset in the DRAM buffer based on the assigned tiles.

```c++
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);

```

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

### Initialize fmod_tile Kernel

Source: https://docs.tenstorrent.com/tt-metal/latest/tt-metalium/tt_metal/apis/kernel_apis/compute/fmod_tile

Initializes the fmod_tile kernel. Refer to the documentation for `any_init` for detailed usage information.

```cpp
void ckernel::fmod_tile_init(uint32_t param0, uint32_t param1)
```

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

### Install Tracy Profiler GUI for Mac using Homebrew

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

Steps for Mac users to install the Tracy GUI using Homebrew. This involves tapping a new formula, downloading the script, and installing it.

```shell
brew uninstall tracy # Remove any old version of Tracy
brew tap-new $USER/tracy # Create a tap
wget -P $(brew --repository)/Library/Taps/$USER/homebrew-tracy/Formula/ --no-check-certificate --no-cache --no-cookies https://raw.githubusercontent.com/tenstorrent-metal/tracy/master/tracy.rb
brew install $USER/tracy/tracy # Install from the tap

# Start the Tracy GUI
tracy

```

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

### Program Setup for Eltwise Binary Operations

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

Initializes the Metalium environment by creating a mesh device, command queue, program object, and defining core coordinates and constants for tile and buffer sizes. This is a foundational step for executing operations on Tenstorrent devices.

```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;

```

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

### Install Tracy Profiler GUI - Mac (Bash)

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

Installs the Tracy profiler GUI on macOS using Homebrew. This involves tapping a new formula, downloading the Tracy Ruby formula, and then installing it.

```bash
brew uninstall tracy # Remove any old version of Tracy
brew tap-new $USER/tracy # Create a tap
wget -P $(brew --repository)/Library/Taps/$USER/homebrew-tracy/Formula/ --no-check-certificate --no-cache --no-cookies https://raw.githubusercontent.com/tenstorrent-metal/tracy/master/tracy.rb
brew install $USER/tracy/tracy # Install from the tap

# Start the Tracy GUI
tracy
```

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

### Configure and Enqueue Kernels in Metalium

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

This snippet demonstrates how to set runtime arguments for reader, writer, and compute kernels in Metalium. It then defines a distributed workload, adds the program to it, and enqueues the workload for execution on the mesh device. Finally, it waits for the computation to finish.

```cpp
SetRuntimeArgs(program, reader, core, {src0_dram_buffer->address(), src1_dram_buffer->address(), n_tiles});
SetRuntimeArgs(program, writer, core, {dst_dram_buffer->address(), n_tiles});
SetRuntimeArgs(program, compute, core, {n_tiles});

distributed::MeshWorkload workload;
distributed::MeshCoordinateRange device_range = distributed::MeshCoordinateRange(mesh_device->shape());
workload.add_program(device_range, std::move(program));
distributed::EnqueueMeshWorkload(cq, workload, false);
distributed::Finish(cq);
```

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

### Example Kernel for Device Profiler

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

An example C++ kernel demonstrating the use of DeviceZoneScopedN within a loop of 'nop' instructions for profiling. This code is part of the 'test_full_buffer' programming example.

```c++
// SPDX-FileCopyrightText: © 2023 Tenstorrent Inc.
//
// SPDX-License-Identifier: Apache-2.0

#include <cstdint>
#include <tools/profiler/kernel_profiler.hpp>

void kernel_main() {
    for (int i = 0; i < LOOP_COUNT; i ++)
    {
        DeviceZoneScopedN("TEST-FULL");
        //Max unroll size
        #pragma GCC unroll 65534
        for (int j = 0 ; j < LOOP_SIZE; j++)
        {
            asm("nop");
        }
    }
}
```

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

### Reader Kernel for Input Data Loading (cpp)

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

This C++ reader kernel is responsible for loading input data from DRAM buffers into circular buffers for computation. It calculates read offsets based on matrix dimensions and the starting tile index for each core. It reads tiles for both input matrices and pushes them to their respective circular buffers. Dependencies include TensorAccessorArgs, TensorAccessor, get_tile_size, and noc_async_read_tile functions.

```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);
            }
        }
    }
}
```
Tokens:
Raw