TT-Metal API 1.0

[eyonland]

---

TT-Metal Overview

TT-Metal is a low-level programming model with user-facing host APIs. This API has not evolved at the pace needed to support the changes within Metal. Therefore, it has become clear that an overhaul is required to support our long-term goals.

---

Current API Limitations

• Inconsistent interfaces and return types
• Clients access internal structures directly (over 100 instances)
• Lack of standardized query/set methods
• Confusing CoreCoord implementation
• Ownership and lifetime of buffers inconsistent

---

Impact

• Difficult for clients to use and maintain
• Increases support and development overhead
• Hinders future scalability and feature additions

---

Schedule for Delivering Work

Overview

As we prepare for release 1.0, stabilizing the APIs is critical to ensure consistency and maintainability across the codebase.

The following changes are required:

• Consistent Type Names: Standardizing type names across the API to improve clarity for users.
• Opaque Types: Ensuring better encapsulation and modularity by making types opaque, preventing external code from directly accessing internal details.
• API Consolidation: Isolating internal objects and their methods, ensuring all interactions occur through well-defined API methods.

The work will progress on an independent branch, metal-api-1.0, with periodic rebasing on the main branch.

---

Tentative Timeline

Time Frame	Risk	Task Summary
October 11, 2024	Low	Introduce inline namespace v0. Begin work on standardizing type names and making types opaque. Initiate new API structure specification.
October 25, 2024	Low	Review and refine new API structure. Prepare for complex API design changes.
November 1, 2024	Medium	Implement new API, consolidate methods, and remove reach-arounds. Conduct internal reviews and stabilize API changes.
November 20, 2024	Low	* Tag stable main with Metal-API-V0.
November 27, 2024	High	* Finalize complex API changes and ensure consistency. Merge changes into the main branch.
December 4, 2024	High	* Complete the deprecation and removal of namespace v0 functionality.

* Denotes impact to branches based on main rather than metal-api-1.0.

---

Work began October 2, 2024.

---

[eyonland]

A Draft PR was done to highlight the boundaries of where we currently perceive our existing api. #13440.

[eyonland]

This PR was merged to main. Note that we did have a few cases that required forward declarations of types within the namespace v0.

[blozano-tt]

I filed a related issue today
Related to #13499

[eyonland]

We would like to invite early feedback on the direction of our proposed changes for Metal API 1.0
These changes are being proposed on the following draft pr #13759

[eyonland]

All, we have recently restructured the declarations for the api. Rather than have the 'api' folder we propose a more conventional approach and moved this folder to tt-metal/tt_metal/include and instead call this tt_metal. This would allow our build process to expose all the includes in this folder as system level includes and the convention from users would be to use #include <tt_metal/metal.hpp>

Additionally, the internal folder has been moved to be under tt-metal/tt_metal/include/tt_metal.

[eyonland]

All, do to resource constraints and prioritization, we have moved our schedule out for "Tag stable main with Metal-API-V0" to November 20th, as well as the subsequent tasks after this. Please let me know if this impacts any dependent schedules.

[eyonland]

All, we have recently made the last backwards compatible change on Metal API V0 and have identified v0.53.0 as the last release. This means that subsequent commits after v0.53.0 risk breaking changes for any libraries external to this repository. Subsequent changes will be against 1.0-beta going forward and this tag will be setup in the near future. The intent is to use the indication of "beta" to mean that we could have breaking changes in Metalium until we stabilize with a new release candidate. At this time we would like to communicate some of the design decisions that will be under review with an upcoming PR whereby the V0 functionality is intended to be removed. The design decisions you see below are subject to change but should give you an opportunity to see and comment. This post will be updated once we finalize the API to reference the final decisions.

Design Decisions for Metal API 1.0

Summary

This document outlines the proposed changes for the upcoming API for the tt-metal library. The goal is to establish a set of design principles and decisions that enhance the API's usability, flexibility, and maintainability.

Design Principles

• Well-Defined Function Contracts: Each function should have clear preconditions (preferably none where reasonable), postconditions, and possible errors. Functions should be "correct by construction," reducing the need for extensive error handling by the user. When feasible, errors should be encapsulated in a return type like std::expected in C++.

• No Mandatory Allocations: Function signatures should not require users to allocate memory to invoke the function or obtain its return value. This promotes efficiency and ease of use.

• Single Responsibility Functions: Functions should perform one action only, adhering to the principle of single responsibility. Documentation should reflect this by listing one action per function, improving clarity and maintainability.

• Opaque Implementation Details: The purpose of each function should be clear without requiring users to understand the implementation details of the opaque objects within the API. This abstraction allows for flexibility and future changes without affecting the user interface.

• Aggregate Types for Optional Parameters: Optional parameters should be encapsulated within aggregate types to avoid the pitfalls of positional defaults. This enhances readability and reduces errors related to parameter ordering.

• Provision of Low-Level Abstractions: The API should provide the lowest level of abstraction necessary, ensuring that other functionalities can be built on top of it. This design choice avoids limiting user capabilities and promotes flexibility.

• Functional API Design: The C++ API should allow for the library to eventually support exposing a C API for easy adoption in other languages. Any limitations to flexibility should be carefully considered to avoid restricting valid use cases.

• Asynchronous Action Handling: All asynchronous actions should provide the ability to query their completion status. This ensures that users can manage and synchronize asynchronous operations effectively.

• Hard requirement on C++17: All design decisions require that some of tt_metal be compiled to be compliant with MISRA standards. At this time we cannot be certain as to which portions of the API will need to be restricted to C++17 so we have limited our work to this standard. This decision is critical for Tenstorrent's mission to support the automotive industry.

1. Opaque Types

Decision Summary

Implement the API using opaque types (handles) instead of transparent types or object methods.

Motivation

• Flexibility: Opaque types allow the library's implementation to change without affecting the API, as handles remain consistent even if underlying resources change. This enables different allocation strategies without exposing implementation details.
• Simplified Object Lifetime Management: Users interact with handles, which are type- and memory-safe, enabling the API to manage resource access and validation.
• Future-Proofing: Allows for alternative implementations or optimizations without impacting the API, facilitating performance improvements and simplifications over time.

Alternatives Considered

1. Transparent Types

   • Cons:
     • Changes to the object structure affect all users, making refactoring and implementation changes difficult without disrupting existing code.

2. Methods Instead of Free Functions

   • Cons:
     • Decentralizes the API, increasing complexity in tracking and documentation.
     • Implies object-centric operations, but creation and some operations would still require static or free functions.
     • Reduces consistency and makes C API parity less feasible.
     • Encapsulating an object-oriented API with a functional one is more cumbersome than the reverse.

Rationale

• Opaque types provide a stable interface for users while allowing internal flexibility.
• Using free functions centralizes the API, enhancing consistency and ease of use.
• Facilitates easier wrapping of the API for C compatibility and higher-level abstractions.

Implications

• Users benefit from a consistent API that shields them from underlying changes.
• The library can evolve internally without forcing users to adapt to changes.

2. SlotMap for Resource Storage

Decision Summary

Use a SlotMap as a centralized, homogeneous container for storing opaque types.

Motivation

• Type and Memory Safety: Ensures handles are valid and correspond to the correct underlying resources.
• Centralized Management: Simplifies resource management and cleanup, reducing the risk of leaks.

Alternatives Considered

1. Raw Heap Allocation

   • Cons:
     • High risk of resource leaks if users forget to release resources.
     • Lacks centralized management for automatic cleanup.

2. Vector

   • Cons:
     • Handles may become invalid due to reallocation.
     • Index-based handles cannot guarantee the resource hasn't changed, leading to potential invalid access.

3. Deque

   • Cons:
     • Additional indirection compared to SlotMap.
     • Less efficient access patterns for certain operations.

4. Shared Pointers

   • Cons:
     • Performance overhead from reference counting
     • Ownership of resources is difficult to track

Rationale

• SlotMap provides stable, type-safe handles and efficient access.
• It avoids the pitfalls of raw pointers and unstable indices, enhancing overall safety.

Implications

• Improved reliability and safety in resource management.
• Slight overhead of SlotMap is justified by the benefits in stability and safety.

3. Use of std::variant for Polymorphism

Decision Summary

Utilize std::variant (e.g., SlotMap<Key, std::variant<FooA, FooB, ...>>) for value-based polymorphism.

Motivation

• Visitor Pattern: Enables clean and efficient polymorphic behavior using visitors.
• By-Value Storage: Avoids extra indirection and heap allocation associated with pointers.
• Performance: Direct calls to concrete types within visitors reduce overhead.
• Compile Errors: SFINAE can be used in C++17 to communicate interface constraints

Alternatives Considered

1. Virtual Polymorphism (e.g., std::vector<std::unique_ptr>)

   • Cons:
     • Multiple indirections and virtual dispatch overhead.
     • Increased heap allocations leading to potential fragmentation.
     • High-level functions require complex overrides or suffer from multiple dispatches.

2. Tuple of Containers (e.g., std::tuple<SlotMap<Key, FooA>, SlotMap<Key, FooB>, ...>)

   • Cons:
     • Increased complexity in managing multiple containers.
     • More data required in handles for type identification.
     • Less efficient and more cumbersome to use.

Rationale

• std::variant offers a balance between performance and flexibility.
• It allows polymorphic behavior without the overhead of virtual functions and heap allocations.
• Simplifies high-level function implementations by avoiding multiple dispatches.

Implications

• Efficient runtime polymorphism with minimal overhead.
• Enhanced performance and simplicity in API usage.

4. Built-in Thread Safety in DevicePool and Similar Components

Decision Summary

Implement built-in thread safety using atomics and wait-free mechanisms for read-only access.

Motivation

• Ease of Use: Relieves users from managing thread safety, reducing complexity and potential errors.
• Performance: Atomics and wait-free implementations offer efficient concurrency control.

Alternatives Considered

1. Decentralized Device Storage

   • Cons: Similar issues as raw heap allocation; increased risk of resource mismanagement.

2. User-Managed Thread Safety

   • Cons:
     • Necessitates locks, which are more expensive than atomics.
     • Adds burden on users, increasing the likelihood of concurrency bugs.

3. Usage Patterns to Circumvent Thread Safety

   • Cons:
     • Difficult to enforce and verify adherence.
     • Incorrect usage can lead to elusive and hard-to-debug concurrency issues.

Rationale

• Providing built-in thread safety enhances reliability and user experience.
• Atomics offer a performant solution suitable for most use cases.
• The benefits outweigh the minimal overhead introduced.

Implications

• Users can safely use the API in multi-threaded contexts without additional effort.
• Reduces potential for concurrency-related bugs.

5. Use of tt::stl::Span for Contiguous Data Passing

Decision Summary

Utilize tt::stl::Span to pass contiguous data without additional allocations.

Motivation

• Efficiency: Allows functions to accept arrays, vectors, and other contiguous data structures without copying.
• Flexibility: Constructible from various contiguous containers like std::vector, std::initializer_list, and C-style arrays.

Alternatives Considered

1. std::vector

   • Cons: Forces allocation and potential copying; less efficient for read-only or external data.

2. Raw Pointers with Size Parameters

   • Cons:
     • Requires separate size parameter, increasing the risk of mismatches and errors.
     • Reduces memory safety and makes function behavior harder to reason about.

3. std::initializer_list

   • Cons: Not constructible from std::vector or arrays; less flexible.

4. Fixed-size Arrays

   • Cons: Restricts functions to specific sizes or requires templating, reducing flexibility.

5. std::span

   • Cons: This was not available for C++17 which is required for tt-metal.

Rationale

• tt::stl::Span provides a non-owning view over contiguous data, combining efficiency with safety.
• It accommodates multiple data sources without forcing allocations or copies.

Implications

• Functions become more flexible and efficient.
• Users can pass data without worrying about unnecessary overhead.

6. Use of tt::stl::AnyRange for Non-Contiguous Ranges

Decision Summary

Implement tt::stl::AnyRange to handle non-contiguous, type-erased ranges in the API, both owning and non-owning.

Motivation

• Abstraction: Allows the API to return ranges without exposing implementation details.
• Flexibility: Supports lazy evaluation and non-contiguous data structures.
• Performance: Avoids unnecessary allocations by enabling lazy ranges for query results.

Alternatives Considered

1. Returning std::vector by Reference

   • Cons:
     • Requires pre-computation or caching, potentially wasting resources.
     • Exposes internal data structures, limiting flexibility.

2. Returning std::vector by Value

   • Cons:
     • Forces eager allocation and initialization, which can be inefficient.

3. Lazy Ranges without Type-Erasure

   • Cons:
     • Leaks implementation details into the API.
     • Reduces the ability to change implementations in the future.

4. Only support Non-owning Type-Erasure (e.g. boost::any_range)

   • Cons:
     • This prevents the Metal implementation of API methods from using lazy views, because while the underlying containers are not owned, the tt::stl::AnyRange must take ownership of the view wrapper in order to avoid a dangling reference when returned.

Rationale

• tt::stl::AnyRange provides a balance between abstraction and performance.
• It allows the API to offer flexible and efficient range-based results without committing to a specific data structure.
• It allows complex queries to be composed using lazy range views rather than performing dynamic allocation to eagerly buffer the results into an owning container that is constructed on-demand.
• The performance cost of type-erasure can be mitigated by constructing a container for faster access, if the use-case is to re-use the range many times and not just once.

Implications

• Enhances the ability to change internal implementations without affecting the API.
• Users benefit from efficient data access patterns.

7. Use of tt::tt_metal::Scoped for RAII handles

Decision Summary

Implement tt::tt_metal::Scoped to provide RAII move-only objects around resource handles that will call the corresponding de-initialization function during destruction.

Motivation

• Convenience: Allows users to opt-in to higher level abstractions with almost no boilerplate.
• Flexibility: Maintains ability for implementation to manage resources without affecting users.

Alternatives Considered

1. Make lowest level use RAII

   • Cons:
     • Makes API inconsistent by implementing behavior with objects rather than functions
     • Handles would either need to be move-only or use std::shared_ptr in order for implementation to be possible.
     • If the API were to provide RAII as a baseline, some existing use-cases would not be supported and require significant refactoring to accommodate.

2. Make API creation functions return std::shared_ptr

   • Cons:
     • Requires resources to be allocated at stable addresses.
     • Encourages less maintainable patterns in user code that restrict future flexibility of implementation.

3. Expect users (or higher-level libraries) to implement their own RAII wrappers

   • Cons:
     • May lead to unnecessary code duplication, or incorrect/inconsistent behavior if wrongly implemented.
     • Would not be able to use any metal library implementation details, which could benefit performance.

Rationale

tt::tt_metal::Scoped implements a centralized and low-boilerplate mechanism for RAII around Metal 1.0 API without precluding the ability to use the lowest level abstraction directly, if desired.

Implications

To obtain unscoped resource handles from creation functions, use auto to deduce the type and manually call the de-initialization function to release:

auto deviceHandle = tt::tt_metal::v1::CreateDevice(...);
auto bufferHandle = tt::tt_metal::v1::CreateBuffer(...);
...
tt::tt_metal::v1::CloseDevice(deviceHandle);
tt::tt_metal::v1::DeallocateBuffer(bufferHandle);

To obtain scoped resource handles from creation functions, use tt::tt_metal::Scoped to deduce the type and allow the destructor to automatically call the de-initialization function:

tt::tt_metal::v1::Scoped deviceHandle = tt::tt_metal::v1::CreateDevice(...);
tt::tt_metal::v1::Scoped bufferHandle = tt::tt_metal::v1::CreateBuffer(...);
...
// bufferHandle destructor will automatically call DeallocateBuffer
// deviceHandle destructor will automatically call CloseDevice

A moved-from scoped handle will not call the de-initialization function when it is destructed, and no longer manages an underlying resource. A move-constructed or move-assigned scoped handle will call the de-initialization function to release the underlying resource unless it is later moved-from:

tt::tt_metal::v1::Scoped deviceHandle = tt::tt_metal::v1::CreateDevice(...);
auto otherDeviceHandle = std::move(deviceHandle);
// deviceHandle no longer manages any underlying resource after this point
...
// otherDeviceHandle destructor will automatically call CloseDevice
// deviceHandle destructor does nothing

Move assignment to a scoped handle that is not yet moved-from will de-initialize its current underlying resource by calling the respective de-initialization function before assuming ownership of the other scoped handle's underlying resource:

tt::tt_metal::v1::Scoped firstHandle = tt::tt_metal::v1::CreateDevice(...);
tt::tt_metal::v1::Scoped secondHandle = tt::tt_metal::v1::CreateDevice(...);
...
// secondHandle will call CloseDevice on its underlying resource
secondHandle = std::move(firstHandle);
// secondHandle now manages the underlying resource obtained from firstHandle

A scoped handle can be "released" to obtain its unscoped handle and opt-out of automatic de-initialization, if desired. This aligns with the design of std::unique_ptr, which also supports the same operation:

tt::tt_metal::v1::Scoped scopedHandle = tt::tt_metal::v1::CreateDevice(...);
auto unscopedHandle = scopedHandle.release();
// equivalent to: auto unscopedHandle = tt::tt_metal::v1::CreateDevice(...);
...
tt::tt_metal::v1::CloseDevice(unscopedHandle);

[marty1885]

I want to comment on the decision on using opaque handles and the piratical impacts it is going to have. Is there any reason to not use a pimpl pattern for device handling? In the vast majority of cases handles are implemented via a table lookup to the real underlying object. Which 1) breaks ASan's ability to track down dangled pointers (no more pointers, just handlers). 2) For large enough tables (ex: buffers) lookup to the real object is going to be (much) slower then a vtable call as the CPU is less likely to have the table in cache. Especially if it's a hash table. 3) Handles are often accidentally castable to each other if not implemented strictly and correctly. Leading to unwanted shooting in own foot, like accidentally passing a device handle into a buffer handle.

I don't see any advantage of an opaque handle compared to a pimpl here. The compiler can do more to help the user with pimpls. And the same advantage of hiding details are kept.

To obtain unscoped resource handles from creation functions, use auto to deduce the type and manually call the de-initialization function to release:

auto deviceHandle = tt::tt_metal::v1::CreateDevice(...);
auto bufferHandle = tt::tt_metal::v1::CreateBuffer(...);
...
tt::tt_metal::v1::CloseDevice(deviceHandle);
tt::tt_metal::v1::DeallocateBuffer(bufferHandle);

This feels like footgun and bugs waiting to happen. Even OpenCL's C++ API does mandatory reference counting for it's buffers (IMO that's a bad design, but it's made for pre-c++11 so I'll give it a pass).

[patrickroberts]

1. Because the handle uses a slot map as the lookup table, even if the resource is freed, the table will be aware when a handle refers to a resource that is no longer available (even if a new resource's lifetime is started at the same index, because the version bits in the key would mismatch), and will throw an error instead of resulting in UB or a segfault.
2. Slot maps are flat, they are not like hash tables, and are more cache friendly, and the expectation is that if the resource is frequently used, the handle should be on the stack or possibly even a register, and the resource's memory would remain in cache.
3. I can assure you there is no valid cast between these handles, and we do not internally cast between them either. They are completely unrelated types as far as the type system goes.

The advantages of an opaque handle are that the library implementation can have control of where the resources are allocated (the purpose of which, as you pointed out, would be to improve cache friendliness for example), and be able to better reason about their lifetime, since it, rather than the user, manages that. Because it doesn't use unique_ptr or shared_ptr, it is also possible for the resource to be relocated without invalidating the handle, in case it is allocated in a container that can reallocate (slot map is implemented using a vector, so it is indeed the case for that).

[nsmithtt]

+1 for using pimpl. I've worked on OpenGL graphics drivers in the past which used integer handles like this and it was a massive pain to debug. e.g. if you break in the debugger, all you have is an integer handle, you'll have to go manually find wherever the slot map is so you can dig up its pointer just to get underlying data out of it. It just adds an indirection for little/no benefit.

It's also not clear to me that slot map is necessarily faster than pimpl for the reasons that Marty gave.

If we're tied to slot map, we could just treat it like an allocator, and still vend a pimpl object, the pimpl object could be a fat pointer that holds both a pointer and a handle.

[marty1885]

Thanks for the replay and addressing the concerns. It sounds like the opaque object idea is only as-good as a pimple if implemented correctly and no additional advantages over pimpl.

The advantages of an opaque handle are that the library implementation can have control of where the resources are allocated (the purpose of which, as you pointed out, would be to improve cache friendliness for example)

You can do that with PMR or custom allocators/deletors too. Then simply wrap that behind a function. The STL supports it as a native functionality and is very well tested. It feels like reinventing the wheel if it's the reason to use opaque handles.

and be able to better reason about their lifetime, since it, rather than the user, manages that.

I argue users should manage their objects. First, users should not pay for features they don't use, especially in libraries like TTNN and Metal where performance is absolutely critical. Second, the library has no business reasoning about the lifetime and how it's used. Users should be the one in control. To say rather then the user managing it, is to make a lot of assumptions on how the library would be used. Which will be wrong from time to time. And be a major pain for the user to work around.

Slot maps are flat, they are not like hash tables, and are more cache friendly, and the expectation is that if the resource is frequently used, the handle should be on the stack or possibly even a register, and the resource's memory would remain in cache.

Feels like a vtable with extra steps. And a pain to debug because the use of integer handles (off by 1 errors, aliasing, garbage data, etc..). The same point as Nick mentioned. Also my point of cache misses is still valid. The CPU would likely prefetched and kept the vtable as it is pointed by the pointer. While the CPU can't for a handle as it sees an integer but nothing to offset into.

[dmakoviichuk-tt]

@patrickroberts

1. Not sure we really need it.
2. It make debug more difficult. Having pointer makes life much easier.
3. Interfaces would have the same properties.

I got a feeling that you are reimplementing OOP and virtual functions which doesn't make a lot of sense to me.
The advantages of an opaque handle are that the library implementation can have control of where the resources are allocated (the purpose of which, as you pointed out, would be to improve cache friendliness for example), and be able to better reason about their lifetime, since it, rather than the user, manages that. Because it doesn't use unique_ptr or shared_ptr, it is also possible for the resource to be relocated without invalidating the handle, in case it is allocated in a container that can reallocate (slot map is implemented using a vector, so it is indeed the case for that).

It is unrelated to the more common c++ approaches. Also reallocating resources at the same slot might generate even funnier bugs for us and I'd better get UB or crash here.
Also if decide to use pimpl it is very easy to make sure pimpl will not require additional memory allocation from heap. Overall I'd prefer to use jsut simple interfaces returned by API. Creation function can return something like std::unique_ptr and everything will be just fine. If I want to have complex ownership I'd just create shared_ptr from it. If single ownership and move only is enough I'll stay with unique.

It feels like you are reinventing a wheel but decided to make it square instead of circle.
I am potential user of your API and I don't like what I see here.

[eyonland]

All, we are halting work on this effort until further notice.