Runtime Arguments

void tt::tt_metal::SetRuntimeArgs(const Program &program, KernelHandle kernel, const std::variant<CoreCoord, CoreRange, CoreRangeSet> &core_spec, ttsl::Span<const uint32_t> runtime_args)

Set runtime args for a kernel that are sent to the core during runtime. This API needs to be called to update the runtime args for the kernel. The number of runtime args per core (unique and common runtime args count toward the same limit) is bounded by the available L1 kernel-config space for the target core type; max_runtime_args is a conservative portable floor.

Return value: void

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

core_spec Location of Tensix core(s) where the runtime args will be written

runtime_args The runtime args to be written

void tt::tt_metal::SetRuntimeArgs(const Program &program, KernelHandle kernel, const std::variant<CoreCoord, CoreRange, CoreRangeSet> &core_spec, std::initializer_list<uint32_t> runtime_args)

Set runtime args for a kernel that are sent to the core during runtime. This API needs to be called to update the runtime args for the kernel. The number of runtime args per core (unique and common runtime args count toward the same limit) is bounded by the available L1 kernel-config space for the target core type; max_runtime_args is a conservative portable floor.

Return value: void

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

core_spec Location of Tensix core(s) where the runtime args will be written

runtime_args The runtime args to be written

void tt::tt_metal::SetRuntimeArgs(const Program &program, KernelHandle kernel, const std::vector<CoreCoord> &core_spec, const std::vector<std::vector<uint32_t>> &runtime_args)

Set multiple runtime arguments of a kernel at once during runtime, each mapping to a specific core. The runtime args for each core may be unique. The number of runtime args per core (unique and common runtime args count toward the same limit) is bounded by the available L1 kernel-config space for the target core type; max_runtime_args is a conservative portable floor.

Return value: void

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

core_spec Location of Tensix core(s) where the runtime args will be written

runtime_args The runtime args to be written

RuntimeArgsData &tt::tt_metal::GetRuntimeArgs(const Program &program, KernelHandle kernel_id, const CoreCoord &logical_core)

Get the runtime args for a kernel.

Return value: uint32_t *

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

logical_core The location of the Tensix core where the runtime args will be written

std::vector<std::vector<RuntimeArgsData>> &tt::tt_metal::GetRuntimeArgs(const Program &program, KernelHandle kernel_id)

Get the runtime args for a kernel.

Return value: std::vector< std::vector< RuntimeArgsData > > &

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

void tt::tt_metal::SetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id, ttsl::Span<const uint32_t> runtime_args)

Set common (shared by all cores) runtime args for a kernel that are sent to all cores during runtime. This API needs to be called to update the common runtime args for the kernel. The number of runtime args per core (unique and common runtime args count toward the same limit) is bounded by the available L1 kernel-config space for the target core type; max_runtime_args is a conservative portable floor.

Return value: void

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

runtime_args The runtime args to be written

void tt::tt_metal::SetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id, std::initializer_list<uint32_t> runtime_args)

Set common (shared by all cores) runtime args for a kernel that are sent to all cores during runtime. This API needs to be called to update the common runtime args for the kernel. The number of runtime args per core (unique and common runtime args count toward the same limit) is bounded by the available L1 kernel-config space for the target core type; max_runtime_args is a conservative portable floor.

Return value: void

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args

runtime_args The runtime args to be written

RuntimeArgsData &tt::tt_metal::GetCommonRuntimeArgs(const Program &program, KernelHandle kernel_id)

Get the common runtime args for a kernel.

Return value: RuntimeArgsData &

Argument Description

program The program containing kernels, circular buffers, semaphores

kernel_id ID of the kernel that will receive the runtime args