Welcome to TT-Metalium!

Get Started

• Getting Started
  • Software Stack Overview
  • Programming Philosophy
  • Installation
  • Key Concepts
  • Quick Start Guide
    • Basic Usage
    • Intermediate Usage
    • Advanced Usage
  • Resources and Debugging
    • Documentation
    • Related Repositories
    • Debugging Tools
    • Community Support
  • Next Steps
• Install
  • Prerequisites:
    • 1: Set Up the Hardware
    • 2: Install Software Dependencies
      • Option 1: TT-Installer Script (recommended)
      • Option 2: Manual Installation
  • TT-NN / TT-Metalium Installation
    • There are four options for installing TT-Metalium:
    • Binaries
      • Step 1. Install the Latest Wheel:
      • Step 2. (For models users only) Set Up Environment for Models:
    • Docker Release Image
    • Source
      • Step 1. Clone the Repository:
      • Step 2. Build the Library:
      • Step 3. Virtual Environment Setup
    • Anaconda
      • Step 1. Install the Latest Package:
  • You are All Set!
    • To verify your installation (for source or wheel installation only), try executing a programming example:
    • Interested in Contributing?
    • Multi-Card Configuration (TT-Topology)
  • Virtual Machine Requirements
    • Overview
    • Why It Matters
    • Requirements for VMs

TT-Metalium

• Programming Model
• Programming Examples
  • DRAM Loopback
    • Device initialization
    • Program setup
    • Create buffers in DRAM and L1 (SRAM)
    • Sending real data into DRAM
    • Creating a data movement kernel
    • Setting runtime arguments for the data movement kernel
    • Running the program
    • Download the result and verify output
    • Validation and teardown
  • Eltwise binary
    • Program setup
    • Circular buffers
    • Data movement and compute kernels
    • Download the result and verify output
    • Validation and teardown
  • Eltwise SFPU
    • Program setup
    • The kernels
    • Set up runtime arguments
    • Program execution and final check
    • Conclusion
  • Matmul (Single Core)
    • Device Initialization & Program Setup
    • Data Preparation and Golden Reference
    • DRAM Buffer Allocation
    • Circular Buffer Orchestration for Pipelined MatMul
    • Matmul Kernel Pipeline Breakdown
      • The reader kernel
      • The compute kernel
      • The writer kernel
    • Kernel execution and result verification
    • Conclusion
  • Matmul (Multi Core)
    • Device Initialization & Program Setup
    • Calculating Work Distribution
    • Buffer and Circular Buffer Allocation
    • Partitioning Work in Kernels
    • Kernel Creation and Parameter Setup
    • Program Execution, Receiving Results and Cleanup
    • Conclusion
    • References
  • Matmul (Multi Core Optimized)
    • Data Reuse in matmul_multicore_reuse
      • Fine-Grained Block Size Control
      • Intermediate Circular Buffer Configuration
      • Stride Kernel Arguments
      • Intermediate Results Handling
      • Conclusion
    • Data Multicasting in matmul_multicore_reuse_mcast
      • Additional Compile-Time Argument
      • Configuring Core Ranges for Tile Distribution
      • Circular Buffer Creation for CoreGrid
      • Multicast Reader/Writer Kernel Setup
      • New Compute Kernel: Fused Bias Addition and Activation Functions
      • Semaphores
      • Kernel Runtime Arguments
  • Writing Custom SFPU Operations using SFPI
    • Vector addition using SFPI
      • Program setup
      • The Kernels
        • Data Movement Kernels
        • SFPI Compute Kernel
      • Custom SFPI Implementation
      • Runtime Arguments and Execution
      • Conclusion
    • Smoothstep using SFPI
      • Building and Running
      • Program setup
      • The Kernels
        • Data Movement Kernels
        • SFPI Compute Kernel
      • Custom SFPI Implementation of Smoothstep
        • Parameter Passing
        • Vector Predicates
      • Runtime Arguments and Execution
      • Conclusion
• Lab Exercises
  • Lab 1: Single Core Matrix Multiplication
    • Introduction
      • Basic Algorithm
    • Row-Major Memory Layout
      • Cache-Friendly Access Patterns
        • Implications for Matrix Multiplication
    • Linear Transformation and Computational Flexibility
    • Loop Tiling
      • Simple Tiling Example
        • Original Doubly Nested Loop
        • Tiled Version
      • Exercise 1: Tiled Matrix Multiplication
    • Introduction to Tenstorrent Architecture
      • Tile-Based Architecture
        • Memory Layout and Tiling
    • TT-Metalium Programming Model
    • Example TT-Metalium Program
      • Exercise 2: Running the Example Program
      • Program Description
        • Kernel Types and Data Flow
        • Kernel Creation and Configuration
        • Reader Kernel Code
        • Compute Kernel Code
        • Writer Kernel Code
      • Example Program Summary
    • Kernel Compilation and Execution
      • Exercise 3: Observing JIT Compile Errors
    • Debug Facilities in TT-Metalium
      • Debug Print API
        • Printing a Tile in a Dataflow Kernel
        • Caveats and Best Practices
      • Exercise 4: Using DPRINT to Debug a Kernel
      • Debugging Hangs using Stack Traces
        • tt-triage Dependencies
      • Exercise 5: Using tt-triage to Debug a Hang
      • Device Performance Profiling
      • Exercise 6: Using Device Profiling to Profile Kernels
    • Matrix Multiplication in TT-Metalium
      • Exercise 7: Implementing Matrix Multiplication in TT-Metalium
    • Conclusion
    • Troubleshooting and Additional Resources
  • Lab 2: Multi Core Matrix Multiplication
    • Introduction
    • From Single Core to Multi Core Matrix Multiplication
      • Work Distribution for Multi Core Programs
      • Work Distribution in TT-Metalium
        • Determine Number of Available Tensix Cores
        • Split Work Among Tensix Cores
        • Create Circular Buffers and Kernels on Multiple Cores
        • Set Per-Core Runtime Arguments
      • Inspecting and Choosing Cores
      • Exercise 1: Multi Core Matrix Multiplication
    • Data Reuse in Multi Core Matrix Multiplication
      • Motivation
      • Blocked Matrix Multiplication
      • Limiting On-Chip Memory Usage
      • Blocked Matrix Multiplication with Data Reuse Pseudocode
      • Data Reuse Evaluation
      • Blocked Matrix Multiplication in TT-Metalium
      • Exercise 2: Multi Core Matrix Multiplication with Data Reuse
    • Potential Additional Optimizations
    • Conclusion
  • Lab 3: Multicast for Improved Data Reuse in Multi Core Matrix Multiplication
    • Introduction
      • High-Level Motivation
    • Background: Tenstorrent NoC and Multicast
    • Example Multicast Program
      • Synchronization with Semaphores
      • High-Level Multicast Protocol
      • Overview of Provided Files
      • Logical vs. Device Coordinates
      • Data Movement Processors and NoC Selection
      • Receiver Kernel Overview
        • Semaphores: Local vs. Remote Access
      • Sender Kernel Overview
        • CB Address Synchronization
        • Multicast Operation
      • Compute and Writer Kernels
      • Multicast and Double Buffering
      • Multicast Exclusions
    • Debugging Hangs with Watcher
    • Exercise 1: Debugging Multicast Issues Using Watcher
    • Exercise 2: Extending the Standalone Multicast Example
    • Exercise 3: Batched Multicast for Improved Throughput
    • Applying Multicast to Multi Core Matrix Multiplication
    • Exercise 4: Multi Core Matrix Multiplication with Multicast
    • Conclusion
    • Appendix A: Watcher Log File Format
• Advanced Topics
  • Hardware implications and the effects
    • Tiles
      • Internal structure of a Tile
      • Conversion between tiles and row-major format
    • Memory from a kernel developer’s perspective
      • Data addressing on Tenstorrent processors
        • RISC-V Address Space
        • DRAM tiles
        • Memory access via the NoC
      • Tensor Layout
      • Memory placement
        • Lock step allocation
        • Interleaved memory
        • SRAM buffers
        • Sharded tensor
    • Compute Engines and Data Flow within Tensix
      • Component introduction
      • Dst register
      • Matrix engine/FPU
      • Vector engine/SFPU
    • Achieving FP32 Accuracy for Computation
      • Host-Side Configuration
      • Kernel-Side Implementation
      • Distinguishing Between matrix and vector engine APIs
• APIs
  • Host APIs
    • Program
      • CreateProgram
    • Buffers
      • CreateBuffer
      • CircularBuffers
      • CreateSemaphore
      • DeallocateBuffer
      • AssignGlobalBufferToProgram
    • Device Management
      • CreateDevice
      • CloseDevice
      • QueryDevices
    • Kernels
      • CreateKernel
      • CreateKernelFromString
    • Runtime Arguments
    • Profiler
      • ReadMeshDeviceProfilerResults
      • GetLatestProgramsPerfData
      • GetAllProgramsPerfData
  • Kernel APIs
    • Common APIs
      • Circular Buffer APIs
        • cb_pages_available_at_front
        • cb_wait_front
        • cb_pages_reservable_at_back
        • cb_reserve_back
        • cb_push_back
        • cb_pop_front
        • get_read_ptr
        • get_write_ptr
      • Kernel Argument APIs
        • get_arg_addr
        • get_arg_val
        • get_common_arg_addr
        • get_common_arg_val
        • get_compile_time_arg_val
    • Data Movement
      • Data Movement APIs
        • Device 2.0 Data Movement API Migration Guide
          • Table of Contents
          • Overview
          • Header Files
          • Key Classes
            • experimental::Noc
            • experimental::CircularBuffer
            • experimental::CoreLocalMem<T>
            • experimental::Semaphore
            • Endpoints
          • Migration Patterns
            • NoC Operations
              • Async Read
              • Async Write
              • Async Write with State (Optimized Repeated Writes)
              • Multicast Write
              • Transaction ID Support
            • Circular Buffer Operations
              • Basic Operations
              • Using CircularBuffer with Noc
              • Selecting Read/Write Pointer for Noc Async Read/Write API
            • Semaphore Operations
              • Local Semaphore Operations
              • Remote Semaphore Operations
              • Multicast Semaphore
            • Memory Access
              • Safe Local Memory Access
              • Struct Access
              • Pointer Arithmetic
          • Complete Migration Examples
            • Example 1: Tile Read Kernel
            • Example 2: Core-to-Core Communication
          • Troubleshooting
            • Common Issues
        • noc_async_read
        • noc_async_read_set_state
        • noc_async_read_with_state
        • noc_async_read_one_packet
        • noc_async_read_one_packet_set_state
        • noc_async_read_one_packet_with_state
        • noc_async_read_page
        • noc_async_read_shard
        • noc_async_write
        • noc_async_write_multicast
        • noc_async_write_multicast_loopback_src
        • noc_async_write_one_packet
        • noc_async_write_one_packet_set_state
        • noc_async_write_one_packet_with_state
        • noc_async_write_multicast_one_packet
        • noc_async_write_page
        • noc_async_write_shard
        • noc_inline_dw_write
        • noc_inline_dw_write_set_state
        • noc_inline_dw_write_with_state
        • noc_async_read_barrier
        • noc_async_write_barrier
        • noc_async_atomic_barrier
        • noc_async_full_barrier
        • noc_async_read_barrier_with_trid
        • noc_async_write_barrier_with_trid
        • noc_async_writes_flushed
        • noc_async_posted_writes_flushed
        • get_semaphore
        • noc_semaphore_set
        • noc_semaphore_set_remote
        • noc_semaphore_set_multicast
        • noc_semaphore_set_multicast_loopback_src
        • noc_semaphore_inc
        • noc_semaphore_wait
        • noc_semaphore_wait_min
        • get_noc_addr
        • get_noc_addr_from_bank_id
        • get_noc_multicast_addr
        • Ordering (NOC, CMD_BUF, VC)
    • Compute
      • Compute APIs
        • Synchronization
          • acquire_dst
          • release_dst
          • reg_api
        • Compute
          • Initialization
            • any_init
            • binary_init_funcs
          • Compute (FPU/matrix engine)
            • add_tiles
            • sub_tiles
            • mul_tiles
            • add_tiles_bcast
            • sub_tiles_bcast
            • mul_tiles_bcast
            • matmul_tiles
            • matmul_block
            • reduce_tile
            • transpose_wh_tile
            • tilize
            • untilize
          • Compute (SFPU/vector engine)
            • Basic arithmetic operations
              • add_binary_tile
              • sub_binary_tile
              • mul_binary_tile
              • div_binary_tile
              • rsub_tile
              • rsub_binary_tile
              • power_tile
              • power_binary_tile
              • square_tile
              • sqrt_tile
              • rsqrt_tile
              • recip_tile
              • abs_tile
              • negative_tile
              • binop_with_scalar_tile
              • tiled_prod_tile
            • Integer operations
              • add_int_tile
              • sub_int_tile
              • mul_int_tile
              • rsub_int_tile
              • gcd_tile
              • lcm_tile
              • remainder_tile
              • fmod_tile
            • Exponential and logarithmic functions
              • exp_tile
              • exp2_tile
              • expm1_tile
              • log_tile
              • log1p_tile
              • log_with_base_tile
              • xlogy_binary_tile
            • Comparison and logical operations
              • unary_eq_tile
              • unary_ne_tile
              • unary_gt_tile
              • unary_ge_tile
              • unary_lt_tile
              • unary_le_tile
              • unary_max_tile
              • unary_min_tile
              • binary_max_tile
              • binary_min_tile
              • ltz_tile
              • lez_tile
              • gtz_tile
              • gez_tile
              • eqz_tile
              • nez_tile
              • logical_not_unary_tile
              • isinf_tile
              • isposinf_tile
              • isneginf_tile
              • isfinite_tile
              • isnan_tile
              • heaviside_tile
            • Bitwise operations
              • bitwise_and_tile
              • bitwise_or_tile
              • bitwise_xor_tile
              • bitwise_not_tile
              • bitwise_and_binary_tile
              • bitwise_or_binary_tile
              • bitwise_xor_binary_tile
              • left_shift_tile
              • right_shift_tile
              • binary_shift_tile
                • Left Shift
                • Right Shift
                • Logical Right Shift
            • Rounding and ceiling functions
              • Rounding operations
              • frac_tile
              • trunc_tile
              • round_tile
              • floor_tile
              • ceil_tile
              • clamp_tile
              • threshold_tile
              • sign_tile
              • signbit_tile
            • Hyperbolic functions
              • asinh_tile
              • atanh_tile
            • Special mathematical functions
              • i0_tile
              • i1_tile
              • cumsum_tile
              • alt_complex_rotate90_tile
            • Initialization and utilities
              • compute_kernel_hw_startup
            • Trigonometric functions
              • erf_tile
              • erfc_tile
              • erfinv_tile
              • tanh_tile
              • tan_tile
              • sin_tile
              • cos_tile
              • asin_tile
              • atan_tile
              • acos_tile
              • acosh_tile
              • selu_tile
            • Activation functions
              • relu_tile
              • relu_max_tile
              • relu_min_tile
              • leaky_relu_tile
              • elu_tile
              • gelu_tile
              • sigmoid_tile
              • celu_tile
              • silu_tile
              • prelu_tile
              • softplus_tile
              • softsign_tile
              • hardsigmoid_tile
              • hardtanh_tile
            • Data manipulation and processing
              • fill_tile
              • identity_tile
              • mask_tile
              • where_tile
              • addcmul_tile
              • reshuffle_rows_tile
              • typecast_tile
              • dropout_tile
              • rand_tile
            • Quantization operations
              • quant_tile
              • requant_tile
              • dequant_tile
      • Packing APIs
        • copy_tile
        • pack_tile
        • pack_untilize_block
        • reconfig_data_format
      • Low Level Kernels
        • Overview
        • Compiler Options/Flags
          • Example
        • Details
          • Namespace
          • User Visible Data Types
            • User Visible Constants
            • User Visible Objects
            • Macros
          • Data Type Details
            • vFloat
            • vInt
            • vUInt
          • Library
          • Immediate Floating Point Values
            • Boolean Operators
            • vBool
            • Assigning and Using Constant Registers
            • Assigning LRegs
        • Miscellaneous
          • Register Pressure Management
          • Optimizer
          • SFPREPLAY
          • Tools
        • Pitfalls/Oddities/Limitations
          • Arrays/Storing to Memory
          • Function Calls
          • Register Spilling
          • Error Messages
          • Limitations
• Environment Variables
  • TT_METAL_HOME
  • TT_METAL_RUNTIME_ROOT
• Tools
  • Kernel Debug Print
    • Overview
    • Enabling
  • Device Debug Print
    • Overview
    • Enabling
  • Watcher
    • Overview
    • Enabling
    • Details
    • gdb Integration
    • Example
    • Asserts
    • Pausing
    • Ring Buffer
    • Stack Usage Measurement
    • Debug Delays
    • Eth Link
    • Circular Buffer Out-of-Bounds
    • HW Exceptions (Quasar only)
  • Tracy Profiler
    • Overview
    • Enabling Tracy
    • GUI
      • Installing for Mac users
      • Building for Linux users
      • Starting the GUI
    • Capturing Profiling Data
      • Capturing via Command Line
    • Profiling Host Code
      • C++
      • Python
        • Python Scripts
        • Pytest Sessions
        • Instrumenting Python Code
        • Signposts in Python Code
        • Line-level Profiling
    • Profiling Device Code
  • Device Program Profiler
    • Enabling Device Profiling
    • Instrumenting Kernel Code
    • Controlling Profiling from Host Code
    • Example Walkthrough: test_full_buffer
    • Analyzing Profiler Output
      • Tracy GUI Integration
    • Limitations
  • Lightweight Kernel Asserts
    • Overview
    • Enabling
  • LLK Asserts
    • Overview
    • Enabling
    • What LLK Asserts Validate
    • Examples from tt-llk Code
    • Location in Codebase
    • CI/CD Integration
    • When to Use LLK Asserts
    • Performance Impact
    • Debugging Failed LLK Asserts
    • See Also
  • Inspector
    • Overview
    • Enabling
    • Extending
      • Examples
      • Using in tt-triage
    • Limitations
  • tt-triage
    • Overview
    • Running
    • Arguments
    • Extending
  • NOC Debug Dump
    • Overview
    • Enabling
    • Example
    • Limitations

Resources

• Support
  • Reporting bugs, feature proposals, or support requests
  • Troubleshooting and debugging tips
  • Community
• Contributing as a developer

Indices and tables

• Index

• Module Index

• Search Page