Part 6. Practice · lesson 42 of 42
42. Every primitive¶
This page lists, in the order of the three places a Komorebi program runs in, every call that the tutorial's programs use, with one line of what it does and the lesson that explains it. It is for looking up, not for reading in order. It relies on all of Parts 1 to 5.
A program is Python on the host, kernels on the Tensix cores, and, when it has one, a thread on the X280 cores. Commands go to the Tensix cores from Python and from a thread. The three tables below follow that order.
Python, on the host¶
| Call | What it does | Lesson |
|---|---|---|
km.Program(name, kernels=folder) |
describes what the cores will run | lesson 4, Hello, core |
km.Session() |
the open card | lesson 4, Hello, core |
session.load(program) |
builds, loads and starts the kernels; returns the run | lesson 4, Hello, core |
run.submit(kind, group=(first, count)) |
sends a command, returns a handle | lesson 4, Hello, core |
run.wait(handle, timeout=s) |
waits for the event of a command; a list of handles waits for all | lesson 4, Hello, core, lesson 9, Many commands |
run.stop() |
ends the kernels | lesson 4, Hello, core |
event.kind, event.microseconds |
how a command ended, and the card's time for it | lesson 5, Commands |
args=[...] |
up to ten numbers in a command | lesson 7, Arguments |
result=True, event.result |
asks for a result word, and reads it | lesson 6, A result |
handlers=header |
names for the handler numbers | lesson 8, Handlers |
session.load(program, deadline=s) |
the card's time limit for a command | lesson 10, When it goes wrong |
km.WaitTimeout, km.CommandTimeout |
the host gave up, the card gave up | lesson 10, When it goes wrong |
session.tensor(array, dtype, "tile") |
puts an array in DRAM as tiles | lesson 15, Tiles and tensors |
tensor.read() |
reads a tensor back as an array, after stop |
lesson 15, Tiles and tensors |
session.read(tensor) |
the raw pages of a tensor | lesson 17, Writing to DRAM |
tensor.write(array) |
writes a tensor, before load |
lesson 22, A whole map |
tensor.accessor_args(), compile_args= |
lets a kernel address a tensor | lesson 16, Reading from DRAM |
km.Buffer(index, format, page_bytes, pages) |
a circular buffer in the L1 of every core | lesson 19, Circular buffers |
compute={...} |
arithmetic settings of the compute kernel | lesson 21, The matrix unit |
exclusive=True |
a command waits for earlier ones that share cores | lesson 25, A port |
Program(...).dynamic(...) |
the handler that receives messages | lesson 28, A routed message, lesson 30, Full buffers |
Graph, run_hostfed |
a chain of commands that the card runs by itself | lesson 31, Streams and graphs |
program.thread(...), program.ports(...) |
a thread on an X280 cluster, and its ports | lesson 32, Hello, thread, lesson 37, A core asks the X280 |
run.thread_result(timeout), run.thread.read(name) |
the thread's result, and a region after the run | lesson 32, Hello, thread, lesson 35, Sharing memory |
Kernels, on the Tensix cores¶
| Call | What it does | Lesson |
|---|---|---|
km::Reader, km::Compute, km::Writer, .run(lambda) |
the loop of a kernel, one pass for each command | lesson 4, Hello, core |
km::handler(command), command.a0 to a9 |
which handler, and the arguments | lesson 8, Handlers, lesson 7, Arguments |
writer.result(value) |
the result word | lesson 6, A result |
km::first, km::count, me() |
where a core is in the command's group | lesson 11, A core group |
km::partition(n, rank, count) |
a core's share of n items |
lesson 12, Partitions |
km::group_of(command, me) |
.rank and .size of a core in the group |
lesson 13, The leader |
reader.reserve, writer.wait_front, cb_push_back, cb_pop_front |
pass tiles through a circular buffer | lesson 19, Circular buffers |
noc_async_read_page, noc_async_write_page and their barriers |
copy a tile from or to DRAM | lesson 16, Reading from DRAM, lesson 17, Writing to DRAM |
add_tiles, tile_regs_acquire to tile_regs_release, pack_tile |
add tiles in the matrix unit's registers | lesson 20, Computing on a tile |
matmul_init, matmul_tiles |
multiply tiles | lesson 21, The matrix unit |
km::l1(address) |
a core's L1 as a tensor Python can read | lesson 24, Why cores talk |
km::Arrival, km::emit, km::arrive |
send a tile to another core, and wait for it | lesson 25, A port |
km::reduce, reduce_arrive, reduce_emit, reduce_result |
add the tiles of many cores into one | lesson 26, Reduce |
km::Broadcast |
one core's tile to many cores | lesson 27, Broadcast |
km::Messages: send, next, forward, consume, done, quota |
messages to any core | lesson 28, A routed message, lesson 29, Forward and consume, lesson 30, Full buffers |
km::X280Port: call, tell |
ask and tell an X280 core | lesson 37, A core asks the X280 |
Threads, on the X280 cores¶
| Call | What it does | Lesson |
|---|---|---|
kmx_main(core), kmx_args_get(), kmx_result(code, words, count) |
the entry, the arguments, the result | lesson 32, Hello, thread |
kmx_submit, kmx_wait, kmx_time |
send a command to the grid and wait for its event | lesson 33, A thread that submits |
kmx_parallel(function, argument) |
a function on the cluster's four cores | lesson 34, Four cores |
kmx_region, kmx_invalidate, kmx_flush |
a region as memory, and keeping the cache in step | lesson 35, Sharing memory |
kmx_cluster, kmx_lead, kmx_serve, kmx_parallel_all |
a function on all twelve cores | lesson 36, Twelve cores |
kmx_port_owner, kmx_port_asked, kmx_port_answer, kmx_port_told, kmx_port_open, kmx_port_run |
serve a Tensix core's asks and tells | lesson 37, A core asks the X280 |
What next¶
The reference pages of this site list every call with all its options, and the examples combine them into larger programs. lesson 40, A whole program and lesson 41, Your own kernel are the way in: start from a program that works, and change one thing.
Previous: lesson 41, Your own kernel.