### 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 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= ./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= ./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 dist(-0.5f, 0.5f); std::vector src0_vec(M * K, 0); // Matrix A (MxK) std::vector 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 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= ./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= ./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= ./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= ./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= ./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= ./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(0); uint32_t src1_addr = get_arg_val(1); uint32_t Mt = get_arg_val(2); uint32_t Kt = get_arg_val(3); uint32_t Nt = get_arg_val(4); uint32_t output_tile_start_id = get_arg_val(5); // Starting tile ID for this core uint32_t num_output_tiles = get_arg_val(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(); 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(0); uint32_t num_tiles = get_arg_val(1); // Number of tiles to write uint32_t start_id = get_arg_val(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= ./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{}); ``` -------------------------------- ### Sample Log Output During Execution Source: https://docs.tenstorrent.com/tt-metal/latest/ttnn/_sources/ttnn/tutorials/tutorials/ttnn_visualizer.md.txt Example log output showing TT-NN configuration loading and initialization. ```bash (python_env) /root/tt-metal$ pytest models/demos/yolov4/tests/pcc/test_ttnn_yolov4.py::test_yolov4[0-pretrained_weight_true-0] 2025-08-01 09:20:51.664 | DEBUG | ttnn::73 - Loading ttnn configuration from /root/tt-metal/vis.setup 2025-08-01 09:20:51.665 | DEBUG | ttnn::83 - Initial ttnn.CONFIG: Config{cache_path=/root/.cache/ttnn,model_cache_path=/root/.cache/ttnn/models,tmp_dir=/tmp/ttnn,enable_model_cache=false, \ enable_fast_runtime_mode=false,throw_exception_on_fallback=false,enable_logging=true,enable_graph_report=false,enable_detailed_buffer_report=true, \ enable_detailed_tensor_report=false,enable_comparison_mode=false,comparison_mode_should_raise_exception=false, \ comparison_mode_pcc=0.9999,root_report_path=generated/ttnn/reports,report_name=ttnn_visualizer_tutorial,4042956046390500517} 2025-08-01 09:20:51.754 | info | SiliconDriver | Opened PCI device 4; KMD version: 1.34.0; API: 1; IOMMU: disabled (pci_device.cpp:197) 2025-08-01 09:20:51.758 | info | Device | Opening user mode device driver (tt_cluster.cpp:192) 2025-08-01 09:20:51.758 | info | SiliconDriver | Opened PCI device 4; KMD version: 1.34.0; API: 1; IOMMU: disabled (pci_device.cpp:197) 2025-08-01 09:20:51.761 | info | SiliconDriver | Opened PCI device 4; KMD version: 1.34.0; API: 1; IOMMU: disabled (pci_device.cpp:197) 2025-08-01 09:20:51.764 | info | SiliconDriver | Harvesting mask for chip 0 is 0x80 (NOC0: 0x80, simulated harvesting mask: 0x0). (cluster.cpp:295) 2025-08-01 09:20:51.776 | info | SiliconDriver | Opened PCI device 4; KMD version: 1.34.0; API: 1; IOMMU: disabled (pci_device.cpp:197) 2025-08-01 09:20:51.836 | info | SiliconDriver | Opening local chip ids/pci ids: {0}/[4] and remote chip ids {} (cluster.cpp:157) ``` -------------------------------- ### 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 ```