CircularBuffers

CBHandle tt::tt_metal::v0::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	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::v0::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::v0::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

void tt::tt_metal::v0::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 index	uint8_t	0 to NUM_CIRCULAR_BUFFERS - 1	Yes
page_size	Updated page size in bytes	uint32_t		Yes

void tt::tt_metal::v0::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

Version 1

CircularBufferHandle tt::tt_metal::v1::CreateCircularBuffer(ProgramHandle &program, const CoreRangeSet &core_spec, const CircularBufferConfig &config)

Creates a Circular Buffer in L1 memory of specified cores and adds it to the program.

Parameters:

program – The program to which the buffer will be added.
core_spec – Specifies the cores where the circular buffer will be configured.
config – Configuration for the circular buffer.

Returns:

CBHandle representing the Circular Buffer ID.