CircularBuffers
-
CBHandle tt::tt_metal::CreateCircularBuffer(Program &program, const std::variant<CoreCoord, CoreRange, CoreRangeSet> &core_spec, const CircularBufferConfig &config)
-
Creates a Circular Buffer (CB) in L1 memory of all cores within core ranges (inclusive) and adds it to the program. There can be a total of NUM_CIRCULAR_BUFFERS (32) circular buffers per core. Circular buffers hold data and have an associated config which indicates usage of the address space. If the config is specified for multiple buffer indices, the circular buffer address space is shared and each buffer index can potentially have a unique view of the shared space.
Circular buffers can be dynamically allocated or program-local allocated. If the config is created with an L1 buffer or sets a globally allocated address it is dynamic and shares the same address space as the L1 buffer. Otherwise, the circular buffer address space is managed by the program. Address space for program-local circular buffers does not persist across programs.
Return value: Circular Buffer ID (uintptr_t)
Argument
Description
Type
Valid Range
Required
program
The program to which buffer will be added to
Program &
Yes
core_spec
Either a single logical core, a range of logical cores or a set of logical core ranges that indicate where the circular buffer will be configured
const std::variant<CoreCoord, CoreRange, CoreRangeSet> &
Yes
config
Config for circular buffer
const CircularBufferConfig &
Yes
-
const CircularBufferConfig &tt::tt_metal::GetCircularBufferConfig(Program &program, CBHandle cb_handle)
-
Gets a reference to the config owned by circular buffer at the given circular buffer ID.
Return value: const CircularBufferConfig &
Argument
Description
Type
Valid Range
Required
program
The program containing the circular buffer
Program &
Yes
cb_handle
ID of the circular buffer, returned by
CreateCircularBuffers
CBHandle (uintptr_t)
Yes
-
void tt::tt_metal::UpdateCircularBufferTotalSize(Program &program, CBHandle cb_handle, uint32_t total_size)
-
Update the total size of the circular buffer at the given circular buffer handle. Updating a program-local circular buffer requires all circular buffers in the program to be reallocated.
Return value: void
Argument
Description
Type
Valid Range
Required
program
The program containing the circular buffer
Program &
Yes
CreateCircularBuffers
| CBHandle (uintptr_t) | | Yes | | | total_size | New size of the circular buffer in bytes | uint32_t | | Yes |
-
void tt::tt_metal::UpdateCircularBufferPageSize(Program &program, CBHandle cb_handle, uint8_t buffer_index, uint32_t page_size)
-
Update the page size at specified
buffer_index
of the circular buffer at the given circular buffer handle.Return value: void
Argument
Description
Type
Valid Range
Required
program
The program containing the circular buffer
Program &
Yes
cb_handle
ID of the circular buffer, returned by
CreateCircularBuffers
CBHandle (uintptr_t)
Yes
buffer_index
Circular buffer index to update page size.
cb_handle
must be a circular buffer that had previously programmed this indexuint8_t
0 to NUM_CIRCULAR_BUFFERS - 1
Yes
page_size
Updated page size in bytes
uint32_t
Yes
-
void tt::tt_metal::UpdateDynamicCircularBufferAddress(Program &program, CBHandle cb_handle, const Buffer &buffer)
-
Update the address of a dynamic circular buffer. Dynamic circular buffers share the same address space as L1 buffers.
Return value: void
Argument
Description
Type
Valid Range
Required
program
The program containing the circular buffer
Program &
Yes
CreateCircularBuffers
| CBHandle (uintptr_t) | | Yes | | | buffer | Dynamically allocated L1 buffer that shares address space of circular buffercb_handle
| const Buffer & | L1 buffer | Yes |