Welcome to TT-Metalium documentation!

Get Started

• Getting Started
  • Quick Start Guide
    • 1. Install and Build
    • 2. Beginner Metalium Usage: DRAM Loopback
    • 3. Beginner Metalium Usage: Eltwise Binary Kernel
    • 4. Beginner Metalium Usage: Single-core Matrix Multiplication Kernel
    • 5. Advanced Metalium Usage: Multi-core Matrix Multiplication Kernel
    • Where to go from here

TT-Metalium

• Programming Model
• APIs
  • Host APIs
    • Program
      • CreateProgram
    • Buffers
      • CreateBuffer
      • CircularBuffers
        • Version 1
      • CreateSemaphore
      • DeallocateBuffer
      • AssignGlobalBufferToProgram
    • CommandQueue
      • EnqueueWriteBuffer
      • EnqueueReadBuffer
      • EnqueueProgram
      • EnqueueRecordEvent
      • EnqueueWaitForEvent
      • EventQuery
      • EventSynchronize
      • BeginTraceCapture
      • EndTraceCapture
      • ReplayTrace
      • ReleaseTrace
      • EnqueueTrace
      • Finish
      • Synchronize
    • Device Management
      • CreateDevice
      • CloseDevice
      • QueryDevices
    • Kernels
      • CreateKernel
      • CreateKernelFromString
    • Runtime Arguments
    • Profiler
      • DumpDeviceProfileResults
  • Kernel APIs
    • Data Movement APIs
      • noc_async_read
      • noc_async_write
      • noc_async_read_barrier
      • noc_async_write_barrier
      • noc_async_write_multicast
      • noc_semaphore_set_multicast
      • noc_semaphore_set
      • noc_semaphore_wait
      • noc_semaphore_inc
      • Ordering (NOC, CMD_BUF, VC)
    • 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
    • Compute APIs
      • copy_tile
      • move_copy_tile
      • acquire_dst
      • release_dst
      • reg_api
      • any_init
      • abs_tile
      • add_tiles
      • sub_tiles
      • mul_tiles
      • add_tiles_bcast
      • sub_tiles_bcast
      • mul_tiles_bcast
      • matmul_tiles
      • matmul_block
      • exp_tile
      • exp2_tile
      • expm1_tile
      • relu_tile
      • relu_max_tile
      • relu_min_tile
      • leaky_relu_tile
      • elu_tile
      • erf_tile
      • erfc_tile
      • erfinv_tile
      • gelu_tile
      • heaviside_tile
      • isinf_tile
      • isposinf_tile
      • isneginf_tile
      • isfinite_tile
      • isnan_tile
      • i0_tile
      • logical_not_unary_tile
      • recip_tile
      • sign_tile
      • sqrt_tile
      • rsqrt_tile
      • sigmoid_tile
      • log_tile
      • log_with_base_tile
      • power_tile
      • rsub_tile
      • signbit_tile
      • square_tile
      • reduce_tile
      • transpose_wh_tile
      • tanh_tile
      • tan_tile
      • sin_tile
      • cos_tile
      • asin_tile
      • atan_tile
      • acos_tile
      • ltz_tile
      • eqz_tile
      • lez_tile
      • gtz_tile
      • gez_tile
      • nez_tile
      • unary_ne_tile
      • unary_gt_tile
      • unary_lt_tile
      • cb_wait_front
      • cb_pop_front
      • cb_reserve_back
      • cb_push_back
      • binary_init_funcs
      • tilize
      • untilize
    • Kernel Argument APIs
      • get_arg_addr
      • get_arg_val
      • get_common_arg_addr
      • get_common_arg_val
      • get_compile_time_arg_val
    • 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
          • Grayskull and Wormhole
          • Wormhole only
        • Immediate Floating Point Values
          • Boolean Operators
          • vBool
          • Assigning and Using Constant Registers
          • Assigning LRegs
      • Miscellaneous
        • Register Pressure Management
        • Optimizer
        • SFPREPLAY
        • Emulation
        • Tools
      • Pitfalls/Oddities/Limitations
        • Arrays/Storing to Memory
        • Function Calls
        • Register Spilling
        • Error Messages
        • Limitations
    • Packing APIs
      • pack_tile
      • reconfig_data_format
• Programming Examples
  • DRAM Loopback
    • Silicon accelerator setup
    • Program pre-compilation setup
    • Building a data movement kernel
    • Create buffers in DRAM and L1
    • Sending real data into DRAM
    • Setting runtime arguments for the data movement kernel
    • Running the program
    • Launch and verify output
    • Validation and teardown
  • Eltwise SFPU
    • Circular buffers for data movement to/from compute engine
    • Compile-time compute kernel arguments
    • Compute kernel declaration and compile-time defines
    • Extra runtime arguments for reader/writer
    • Conclusion
  • Eltwise binary
    • New buffers
    • Compute kernel declaration and compile-time defines
    • Extra source tensor
    • Conclusion
  • Matmul (Single Core)
    • Host Code
    • Main blocks in matmul_single_core function
    • Create Program, Enqueue initialization, and core range definition
    • Create DRAM buffers & Circular buffers
    • Compile-time kernels arguments
    • Compute kernel declaration and compile-time defines
    • Runtime arguments and program launch
    • Conclusion
  • Matmul (Multi Core)
    • Accessing all the cores
    • Splitting the work across cores
    • Using different kernels for reader/writer
    • Compute kernel args
    • Reader/writer kernel runtime args
    • Conclusion
  • 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
• Tools
  • Kernel Debug Print
    • Overview
    • Enabling
  • Watcher
    • Overview
    • Enabling
    • Details
    • gdb Integration
    • Example
    • Asserts
    • Pausing
    • Ring Buffer
    • Stack Usage Measurement
    • Debug Delays
  • Tracy Profiler
    • Overview
    • GUI
    • Profiling Host
      • C++
      • Python
        • Python Scripts
        • Pytest Sessions
        • Instrumenting Python Code
        • Signposts in Python Code
        • Line-level Profiling
    • Profiling Device
  • Device Program Profiler
    • Overview
      • How to Run
    • Example
      • Description
    • 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