Skip to content

API reference

A thorough, type-annotated reference for the public API. For a one-screen scan of the whole surface, see the Cheat sheet; for the raw Doxygen-extracted signatures, see Autodoc.

Everything is in namespace tny (namespace cs = cuda::std). Include <teeny/teeny.h> for all of it; <teeny/cuda.h> adds the CUDA memory spaces.

Legend: Idx = the shape's index_type (int64_t for shape<...>). T = element type. A static index is an integral_constant (Int<k>()); a runtime index is a plain integer. "→ view" means a non-owning tensor<…, storage::view> aliasing the same memory (no copy).


The tensor type

template <class T, class Shape, class Layout = ccontiguous, storage O = storage::view>
struct tensor;
Parameter Meaning Typical value
T element type any arithmetic type, half, bfloat16
Shape the shape (exposed as shape_type / extents_type) shape<2,3> (a cs::extents<int64_t,…>; -1 = dynamic)
Layout memory order ccontiguous (C, default), fcontiguous (F), dynamic_strides (runtime), strides<S...> (static/mixed)
O ownership storage::view (default), storage::stack, storage::heap, storage::gpu/pinned/mapped, storage::gpu_view

Slicing / permuting / peeling / .at() of a gpu tensor yields a storage::gpu_view (a non-owning view of device memory), so a device pointer is never mistaken for a host one in the type. Helpers: storage_is_device (gpu/gpu_view), storage_is_host_accessible, storage_is_view (view/gpu_view/pinned_view/mapped_view), storage_view_of(O) (the view kind that preserves a source's space). pinned/mapped are host-accessible, but a view of one keeps its space — pinned_view / mapped_view (so DLPack can label it kDLCUDAHost) — while behaving like a plain view otherwise.

Ownership aliases

Alias Ownership Notes
view<T,E,L=ccontiguous> none (host view) trivially copyable, kernel-passable; the bare tensor is this
local<T,E,L> stack requires a fully static shape; sizeof == its data
owned<T,E,L> heap (host) move-only
gpu<T,E,L> / pinned<T,E,L> / mapped<T,E,L> CUDA from <teeny/cuda.h>; a view of a gpu is storage::gpu_view

Factories

Wrap existing memory (→ view):

Call Returns Notes
wrap(ptr, shape) view<T,E> C-order (ccontiguous)
wrap<Layout>(ptr, shape) view<T,E,Layout> chosen layout
wrap(ptr, shape, {s0,s1,…}) view<T,E,dynamic_strides> runtime strides (elements; may be negative). A stride-0 axis self-overlaps — fine to read (broadcast), but an in-place write is a host-debug error (clone() first)
wrap<S...>(ptr, shape, {dyn…}) view<T,E,strides<S...>> mixed static/runtime strides (dynamic_stride slots)
wrap(ptr, shape, strides<S...>{}) view<T,E,strides<S...>> compile-time strides (fold into the type)
wrap(…, storage_v<storage::gpu>) view in that space trailing memory-space tag on any overload above (default storage::view); pass the plain backend — storage::gpu/pinned/mapped — and it folds to the view kind (gpu_view/…), since wrap always views. storage_c<S>{} / storage_v<S> are the braced / no-braces spellings
as_tensor(md) view<…> wrap any cs::mdspan/submdspan result
make_view(ptr, shape) view<T,E> an alias of wrap that deduces E (make_view<Layout> for the layout); takes the same trailing storage_c<Space>{} tag

Input → output type — wrap view facets. The element type T is deduced from the pointer (a const T* gives a read-only view); the extents come from the passed shape (each axis static/dynamic exactly as in its E); the ownership is always a view. Only the layout and the space vary with the call form (storage_view_of folds a named backend to its view kind, so wrap never yields an owner):

Argument form Output layout Output space
wrap(ptr, e) ccontiguous storage::view
wrap<Layout>(ptr, e) Layout (e.g. fcontiguous) storage::view
wrap(ptr, e, {s0,s1,…}) dynamic_strides (all runtime) storage::view
wrap<S...>(ptr, e, {dyn…}) strides<S...> (mixed; dynamic_stride slots runtime) storage::view
wrap(ptr, e, strides<S...>{}) strides<S...> (all compile-time, folded) storage::view
wrap(…, storage_v<storage::gpu>) any of the above storage::gpu_view
wrap(…, storage_v<storage::pinned>) / …<storage::mapped> any of the above pinned_view / mapped_view

Allocate new storage — element type T defaults to float; static shape → stack (host+device), dynamic shape → heap (host only):

Call Returns Element type
empty<T>(shape) local/owned (deduced) T (=float); UNINITIALISED
empty<T, storage::S>(shape) owner in space S name a backend: stack/heap/gpu/pinned/mapped
empty<T>(shape, storage_c<storage::S>{}) owner in space S value-tag backend form (same result)
make_local<T>(shape) local<T,E> T (=float); = empty<T,storage::stack>
make_heap<T>(shape) owned<T,E> T (=float); = empty<T,storage::heap>
make_gpu<T>(shape) / make_pinned<T> / make_mapped<T> CUDA owner T (=float); = empty<T,storage::gpu/…>
zeros<T>(shape) / ones<T>(shape) stack or heap T (=float)
full(shape, v) stack or heap the value's type (full<T>(…) to force)
arange<T>(n) owned<T, shape<-1>> T (=int64_t); 1-D [0,n)
arange<T,N>() / arange<T>(Int<N>()) local<T, shape<N>> static 1-D [0,N)
zeros<T, storage::S>(shape) (also ones/full/arange) owner in space S host-accessible backend (stack/heap/pinned/mapped); storage::gpu static_asserts → to<storage::gpu>(zeros<T>(shape)). Value-tag: zeros<T>(shape, storage_c<storage::S>{})

Geometry

Call Returns Notes
t.rank() size_t (constexpr) number of axes
t.numel() integral_constant if fully static, else Idx number of elements (product of the shape); folds when static
t.shape() array-like accessor the tensor's shape (numpy/pytorch spelling): shape()[Int<k>()] folds (integral_constant), shape()[i] is runtime; has rank(), is iterable, converts to the raw Extents
t.shape(d) integral_constant/Idx size of axis d — static Int<k>() folds, runtime d is a value
t.strides() array-like accessor the tensor's strides: twin of shape()strides()[Int<k>()] folds where derivable, strides()[i] is runtime
t.stride(d) integral_constant if derivable, else Idx stride of axis d; static-stride / contiguous layouts fold
t.is_contiguous() bool C-order (numpy/pytorch default) — what reshape/flatten need; is_contiguous(fcontiguous{}) for F
t.is_dense() bool dense block in some axis order (C, F, or permuted); is_dense(L{}) is the exact-layout check and is_contiguous() == is_dense(ccontiguous{})
t.data() T* base pointer
t.view() view<T,E,L> (gpu_view if device) non-owning teeny view aliasing t's memory (no copy)
t.extent(d) / t.extent(Int<k>()) Idx / integral_constant if static mdspan-side per-axis size (== t.shape(d)); folds when static
t.extents() / t.mapping() const Extents& / const mapping& the raw mdspan objects (interop escape hatch)
t.mdspan() cs::mdspan<T,E,L> the raw cuda::std::mdspan

mdspan equivalent

t.shape() and t.strides() are teeny's array-like accessors — indexable with a static Int<k>() (folds to an integral_constant) or a runtime integer, iterable, and rank-aware — and are the primary spelling. t.extents() and t.mapping() hand back the raw cuda::std objects (cs::extents, the layout mapping) for mdspan interop. Per-axis, t.extent(d) == t.shape(d), and t.stride(d) is the strides twin. Reach for extent/extents/mapping/ mdspan() only when you need the underlying mdspan types.


Indexing & slicing

Axis args accept a runtime integer, a static Int<k>(), or a slice specifier. Negative integer indices wrap (count from the back).

Call Returns Notes
t(i, j, k) (all integers) T& element access
t.at(i, j, k) (all integers) rank-0 view scalar-like: converts to/from T, .item(), has atomic_add_
t(0, all, slice(1,4)) (any slice arg) → view lower-/same-rank
t(1, ellipsis, 2) → view or T& ellipsis = rank − #other args copies of all (≤1 per call)
t.take_along<Axes...>(args...) → view bind named axes only, keep the rest
t.take_along(axis<Axes...>{}, args...) → view value form — axis<...> selector (numpy-like), no .template on a dependent receiver

Slice specifiers:

Spelling Meaning
all keep the whole axis (== slice(none,none); folds, keeps static extent)
slice(a, b) / slice(a, b, step) half-open [a,b), optional (negative) step
slice<a,b>() / slice<a,b,step>() compile-time bounds (fold like all)
none open slice end (python None)
ellipsis fill the middle with all

Assignment into a slice copies (broadcasts); on a named view it rebinds:

Statement Effect
a = b (named view) rebind a to b's memory (shallow — nothing copied)
a(ellipsis) = b / a(0,all) = b copy b's elements into the region (broadcasts)
a(...) = 5.0 fill the region

Input → output type — operator() per-axis arg. Element type is unchanged (const if the source is), the output layout is a folded strides<...>, and the space is storage_view_of(O). Each axis argument independently decides whether that axis survives and, if so, its output extent/stride staticity. The output rank = source rank − (number of integer args); ellipsis expands to rank − #other args copies of all before this table applies:

Axis arg Axis in output? Output extent Output stride
i (runtime int) dropped (folded into base offset)
Int<k>() (static int) dropped
all (== slice(none,none)) kept the source extent (static stays static) source stride, folded where derivable
slice<a,b>() / slice<a,b,step>() (static bounds) kept static when the source extent is static (b−a, ceil-div by step) folded (source_static_stride × step) where derivable
slice(a,b) / slice(a,b,step) (runtime bounds) kept dynamic (-1) folded where derivable, else runtime

So t(i, all, slice<1,4>()) on <A,B,C><B,3>; t(i, all, slice(a,b)) on <A,B,C><B,-1>. A negative step (slice(none,none,-1)) reverses the axis (a negative folded stride); a strides<...> source stays fully sliceable.


Structure (views)

All return a view and work on any source layout (incl. strides<...>); axis template args are signed (negatives count from the back). Each <Ax> op also has a value form — pass Int<Ax>() in place of the template argument (t.squeeze(Int<1>()) == t.squeeze<1>(), likewise permute/flip/unsqueeze/ reshape/recast); the value spelling is the preferred one. See Views & structure.

Type inference — what the output type keeps. A view op transforms the input type along four independent facets; staticity is preserved wherever it is derivable, so a kernel written against a static shape stays static through these ops (no accidental fallback to runtime extents/strides):

Facet Rule
extents each kept axis keeps its source extent — static stays static. Peeling the batch dims off shape<-1,-1,M,N> yields a <M,N> cell; a static-bounds slice/range keeps a static extent, a runtime range is dynamic; unsqueeze inserts a static 1
strides folded to a compile-time strides<...> slot wherever source_stride × step is known at compile time — a static source keeps folded strides, and a partially-dynamic contiguous stride folds from the static extents it spans (shape<-1,3,3>stride0 = 9); otherwise runtime
layout the view ops (operator(), take_along, permute, flip, un/squeeze, peel) output a folded strides<...>. recast<keep_strides> (the default) instead preserves the source layout type (ccontiguous stays ccontiguous, strides<> stays strides<>)
space storage_view_of(source): a gpu/gpu_view source → gpu_view, pinned/mapped → the matching _view, else view — a view never loses its memory space

Worked input→output shapes (E = source extents):

Call E → output
t(i, all, slice<1,4>()) <A,B,C> <B,3> — integer drops the axis, all keeps B, a static range keeps a static extent
t(i, all, slice(a,b)) <A,B,C> <B,-1> — a runtime range's extent is dynamic
peel_front<N>(t) cell <*batch(N),M,N> <M,N> — trailing extents and strides stay static even when the batch is dynamic
t.recast<NewE>() <-1,M,N> (any layout) NewE with the static dims recovered; strides + layout preserved
sum<Ax>(t) <A,B,C> <A,C> — named axis removed (static result → stack, any dynamic → heap)
Call Returns Notes
t.permute<Perm...>() → view reorder axes (a permutation of 0..N-1)
t.flip<Ax>() → view reverse an axis (negative-stride)
t.unsqueeze<Ax>() → view, rank+1 insert a size-1 axis
t.squeeze<Ax>() / t.squeeze() → view, rank−1 / − all size-1 drop size-1 axis / axes
t.reshape<NewExt...>() → view contiguous reshape (one -1 inferred)
t.flatten() → 1-D view ravel; needs C-contiguous
t.recast<NewShape[, NewLayout]>() → view reinterpret with a more-static same-rank extents; NewLayout defaults to keep_strides (preserve the source strides AND layout type, any layout, no copy). ccontiguous/fcontiguous = reinterpret AS that order (derive+fold the strides — the "I promise it's contiguous" form; a debug build verifies the imposed strides match the source's and aborts a false promise, symmetric with the extent check — UB only under -DNDEBUG); strides<S...> = impose them. Functional form t.recast(shape{…}, layout{…})
t.reindex<Idx2>() (free: reindex<Idx2>(t)) → view no-copy, layout-preserving retype of the offset index width to Idx2: same pointer + layout kind, extents & dynamic strides narrowed to Idx2 (a strides<...> literal pack unchanged). Narrows a boundary view to shape32 — halves the by-value footprint, 32-bit device offset math. Orthogonal to recast (which staticizes the extent values); they compose. Debug-checks index_fits<Idx2>(); UB if the caller lies
t.index_fits<Idx2>() (free: index_fits<Idx2>(t)) bool do all element offsets fit Idx2? Signed reach Σ_{s>0}(e−1)s / Σ_{s<0}(e−1)sIdx2 (handles negative-stride/flipped views; broadcast stride 0 adds 0)
t.clone() owning (stack/heap) materialise a dense row-major copy. Copies on the HOST → the dynamic-shape overload static_asserts the source is host-accessible; for a gpu/gpu_view tensor use the free to<Space>(t) below
t.to<T2>() view (no-copy) or owning dtype convert. Matching dtype (no Force) → a read-only borrow (gpu_view if t is on the device, else view); differing dtype or t.to<T2,true>() → a dense owning copy (static→stack, dyn→heap). The dynamic-shape copy runs on the HOST and static_asserts host-accessibility — convert a gpu/gpu_view tensor via the free to<Space>(t)
to<Space>(t) (cuda.h) view (no-copy) or owning memory-space move: to<Space, ET, Force>(t)Spacestorage::gpu/pinned/mapped/heap/stack. Same no-copy/Force rule; a device source (gpu/gpu_view) downloads via cudaMemcpy. rvalue source → always copies

reshape/flatten need exact C-contiguity (is_contiguous()); clone() first if the source isn't. recast does not — it only re-types the extents and keeps the source strides, so it works on any layout.

Input → output type — recast<NewE[, NewL]>(). No copy in any form; the element type T and the space are preserved (the result is storage_view_of(O)), the rank is unchanged, and every static dim of NewE recovers (folds into the type) a possibly-dynamic source dim — e.g. <-1,-1>.recast<shape<-1,3,3>>() gives <-1,3,3>. The NewL argument chooses the strides + layout type:

NewL Output extents Output strides Output layout type
keep_strides (default) NewE the source strides, unchanged the source L (ccontiguousccontiguous, strides<S...>strides<S...>, …)
ccontiguous / fcontiguous NewE derived + folded from NewE (C / F packing) ccontiguous / fcontiguous
strides<S...> NewE the imposed S... (dynamic slots filled from the source) strides<S...>

An explicit NewL is a reinterpretation ("I promise it's laid out this way"); a debug build verifies the imposed/derived strides against the source's actual strides (symmetric with the static-extent check), UB only under -DNDEBUG.

Input → output type — to / clone. Whether the result is a borrow or an owning copy, and its ownership, per call:

Call Copy? Output type Notes
t.to<>() (same dtype, no Force) no — borrow view of const T (gpu_view if device), source L kept read-only alias, no allocation
t.to<T2>() (differing dtype) yes static → local<T2,E>; dynamic → owned<T2,E> (host) dense ccontiguous copy
t.to<T2,true>() (Force) yes as above (even when T2==T) force-materialise
t.clone() yes static → local<T,E>; dynamic → owned<T,E> (host) dense row-major (ccontiguous) copy
to<Space>(t) (cuda.h) no-copy if already in Space, else yes borrow (view-kind of Space) or owner in Space memory-space move; Spacegpu/pinned/mapped/heap/stack; rvalue source always copies

The to<T2>()/clone() copying paths run copy_ on the HOST, so they cannot dereference DEVICE memory; the dynamic-shape (_TNY_HOST) overloads static_assert the source is host-accessible. To copy or convert a gpu/gpu_view tensor use the free to<Space>(t) (cuda.h) above, which is device-aware.

nd-peel (iteration)

Call Returns Notes
peel<Axes...>(t) a range of views iterate the named axes; each item is a lower-rank view. The range-for is incremental — it advances the base pointer and reuses the loop-invariant cell mapping (O(1) per step), not an O(#peeled-axes) decode per cell
peel<Axes...>(t).subrange(lo,hi) a [lo,hi) sub-range seed the incremental cursor once at lo, then O(1)/step — split [0,size()) across CPU threads / device blocks, each sweeping its own contiguous chunk
peel<Axes...>(t).enumerate() a range of {index, cell} range-for that ALSO yields the peeled multi-index: for (auto [m, cell] : …) with m[d] = the coordinate of peeled axis d (row-major, last listed fastest) — for a per-axis table axtab[d][m[d]] or a write-by-coord. Opt-in (the bare peel cell stays lean — no coordinate words); composes with subrange (.enumerate().subrange(lo,hi)). The raw iterator also exposes it.index(d) / it.index()
peel_at<Axes...>(t, i) → view the i-th sub-view — random access, decoded from scratch; this is what a device grid-stride loop (i += nthreads) needs
peel_front<N>(t) a range of views N≥0: peel the first N axes; N<0: keep the last |N|. Same incremental range-for + subrange(lo,hi)
peel_front_at<N>(t, i) → view the i-th (random access / grid-stride)
size_front<N>(t) → index # cells peel_front<N> yields (product of the peeled extents), no range built

Input → output type — peel cell. Each yielded cell is a view (storage_view_of(O)gpu/gpu_view source → gpu_view), element type T unchanged, layout a folded strides<...>. The kept axes carry their source extents (static stays static) and folded strides; a peel_front<N> cell over <*batch(N), M, N> is <M, N> with the trailing extents and strides static even when the batch is dynamic. peel<Axes...> drops the named axes; the cell rank is source rank − #peeled axes.


Math

Element type of results follows promote(A,B) (C++ rules, but among floats the lower width wins — pytorch-style; -DTNY_STD_PROMOTION opts out). The result's offset index type is the wider of the two operands' (a mixed int32/int64 broadcast → int64), independent of the element promotion.

In-place (mutates *this, returns tensor&)

A tensor rhs broadcasts (numpy-style; needs equal rank — unsqueeze first); a scalar rhs applies to every element.

Call Notes
a.add_(x) a.sub_(x) a.mul_(x) a.div_(x) x = tensor (broadcasts) or scalar
y.add_(x, alpha) y.sub_(x, alpha) fused scaled accumulate (BLAS axpy): y ± alpha*x, x broadcasts; scaled copy y = alpha*x is y.zero_().add_(x, alpha)
a.atomic_add_(x) a.atomic_sub_(x) atomic accumulate on device (x = tensor or scalar); underlying form is add_<Atomic>/sub_<Atomic>
a += x a -= x a *= x a /= x compound-assign sugar
++a --a prefix, in place
a++ postfix → pre-value stack copy (static shape only)
a.neg_/abs_/exp_/log_/sin_/cos_/sqrt_/tanh_() unary, in place
a.floor_/ceil_/round_/trunc_/sign_() a.pow_(e) a.clamp_(lo,hi) unary, in place
a & b a \| b a ^ b ~a a &= b bitwise (integer element types only)
a.fill_(v) a.zero_() a.copy_(b) a.iota_(start,step) assignment / init (b broadcasts)
a.map_(f) a.zip_with_(g, b) user functor (a device-safe struct, not a lambda)
a.at(i...).atomic_add_(v) scatter-accumulate a(i...) += v (atomic on device)

Out-of-place (→ new tensor; static shape → stack, else heap)

Call Returns
a + b / a.add(b), a - b, a * b, a / b new tensor (tensor+tensor broadcasts, or +scalar)
2.0 * a, 2.0 - a, 1.0 / a, -a scalar–tensor / unary minus
a.pow(b) element-wise power
neg/abs/exp/log/sin/cos/sqrt/tanh/floor/ceil/round/trunc/sign(a) unary free functions — also methods (a.exp(), …)
minimum(a,b) maximum(a,s) clamp(a,lo,hi) elementwise min/max/clamp — also methods (a.minimum(b), …)
normalize(a) cross(a,b) unit vector / 3D cross — also methods (a.normalize(), a.cross(b))
a.add(b, alpha) a.sub(b, alpha) fused out-of-place axpy: a ± alpha*b (b broadcasts); the in-place twin is add_(b, alpha)
a.map(f) new tensor from a user functor

into(dest) — write into a preallocated buffer. Pass into(y) as the last argument to any producer above to write the result into y (one fused pass, no allocation) and return y&, instead of a fresh tensor — the kernel-friendly out=. into(y) is a distinct type, so it never collides with a scalar argument (which is what lets add/sub also take the fused alpha). y may alias an operand and may have a different dtype (the result is cast to it); its extents are checked.

Call Returns
a.add(b, into(y)) a.mul(b, into(y)) a.add(s, into(y)) y& (elementwise, tensor or scalar rhs)
a.add(b, alpha, into(y)) a.sub(b, alpha, into(y)) y& (fused axpy into y)
exp(a, into(y)) … (every unary) y&
minimum(a,b,into(y)) maximum(a,s,into(y)) clamp(a,lo,hi,into(y)) y&
normalize(a, into(y)) cross(a,b,into(y)) y&
sum(a, into(cell))dot(a,b,into(cell)) cell& (full reduction → rank-0 dest)
sum<0>(m, into(buf)) mean(m, axis<1>{}, into(buf)) buf& (axis reduction → lower-rank dest)

Comparisons → a bool tensor (broadcast), reduced with .all()/.any()

Call Returns
a < b, a == 2.0, 3.0 < a, … (== != < <= > >=) bool tensor
(a > 0).all() / (a > 3).any() bool (members)

Reductions

Accumulate in the reduce typedouble for small floats (float/double/half; a wider float keeps its own type), the item type for integers — then cast the result back to the tensor's element type T (sum(float_tensor)float, computed in double). A leading type argument makes that type both the accumulator and the result.

Every reduction here (plus sqnorm/norm below) is also a methoda.sum(), a.mean(), a.dot(b), a.sum<0>(), a.mean(axis<1>{}), a.norm(), a.sum(into(cell)) — with the same overload shapes as the free form. The free sum(a) spelling stays (generic code and the numpy reader both use it).

Call Returns Notes
sum(a) prod(a) max(a) min(a) mean(a) T (accumulated wide) over all axes
dot(a, b) promote(Ta,Tb) (accumulated wide) inner product; extents must match exactly (no broadcast)
sum<Acc>(a), mean<Acc>(a), dot<Acc>(a,b) Acc force the accumulator/return type
allclose(a, b, rtol=1e-5, atol=1e-8) bool \|a−b\| ≤ atol+rtol·\|b\| everywhere (broadcasts)
sum<Axes...>(a) mean<Axes...> max/min/prod<Axes...> lower-rank tensor remove the named axes (negatives wrap)
sum<Acc, Axes...>(a) lower-rank tensor leading type = accumulator, leading int = axis
sum(a, axis<Axes...>{}) (also mean/max/min/prod) lower-rank tensor value form — numpy axis= selector; sum<Acc>(a, axis<...>{}) for the accumulator; no .template on a dependent receiver

Axis reductions: a fully static result → stack (host+device); any dynamic result → heap (host only).

Reductions also take into(dest) (the out= spelling). A full reduction writes its scalar into a rank-0 destination — sum(a, into(cell)), dot(a, b, into(cell)), where cell is a local<T, shape<>>{} or a rank-0 view over an address wrap(&x, shape<>{}) (a non-rank-0 dest is a static_assert, so forgetting the <axes> fails to compile instead of splatting the total). An axis reduction copies its lower-rank result into destsum<0>(m, into(buf)), and the axis<...> value form takes into too (mean(m, axis<1>{}, into(buf))); its dest is broadcast-compatible with the reduced shape (it goes through copy_). Dtype casts, and the destination is returned by reference.

Vector algebra & geometry

Contained exact linear-algebra / geometry on views (no solves, inversion, or optimisation — those stay out of teeny).

Call Returns Notes
sqnorm(a) T (accumulated wide) Σ aᵢ² over all axes; == dot(a, a). sqnorm<Acc> forces accumulator+result
sqnorm<Axes...>(a) / sqnorm(a, axis<...>{}) lower-rank tensor Σ aᵢ² over the named axes (reduction API, like sum); sqnorm<Acc,Axes...>
norm(a) floating (T for float T, double for integer T) √Σ aᵢ² (L2 / Frobenius over all axes); norm<Acc> forces accumulator+result
norm<Axes...>(a) / norm(a, axis<...>{}) lower-rank tensor (floating) √Σ aᵢ² over the named axes; norm<Acc,Axes...>
a.normalize_() tensor& in place a /= norm(a); floating element types only; zero vector → NaN
a.normalize_<Axes...>() tensor& in place, over the named axes (keepdim broadcast); axes distinct & ascending
normalize(a) new tensor (floating; integer → double) out-of-place unit vector; static → stack, dynamic → heap
normalize<Axes...>(a) / normalize(a, axis<...>{}) new tensor (floating) out-of-place, over the named axes (keepdim)
cross(a, b) new stack 3-vector promote(Ta,Tb) 3D cross product; operands rank-1, length 3
a.cross_(b) tensor& in place: a becomes a × b (rank-1, length 3; aliasing-safe). Into a separate slot: slot.copy_(cross(a,b))

Input → output type — result element type. The accumulator and the result dtype are separate. Default accumulator = reduce_type_t<T> (double for a small float float/double/half/bfloat16; a wider float keeps itself; int64_t/uint64_t for a narrow integer, so sum/prod/dot can't overflow mid-accumulation; an already-≥8-byte integer keeps itself). A leading type arg makes that type both the accumulator and the result.

Call Default result dtype With leading Acc
sum(a) prod(a) max(a) min(a) T (item type — accumulate wide, cast back; sum(int8)int8) Acc (e.g. sum<int64_t>(a) keeps the untruncated wide value)
mean(a) T for a floating T; double for an integer T (numpy: integer mean is float64, divided in double) Acc
dot(a, b) promote(Ta,Tb) Acc
sum<Axes...>(a) prod/max/min<Axes...> T (lower-rank tensor) Acc
mean<Axes...>(a) T for floating T; double for integer T (lower-rank tensor) Acc

The axis forms follow the same element-type rule; only the shape changes (named axes removed) and the ownership splits static → stack / dynamic → heap.


Half precision

Name Meaning
half / bfloat16 IEEE binary16 / bfloat16 (native __half/__nv_bfloat16 under nvcc)
compute_type<T> / compute_type_t<T> float for half types, else T (math computes here)

Dispatch (the anyrank boundary)

Call Returns Notes
as_anyrank(data, shape, stride, ndim) anyrank (view store) wraps the arrays, no copy (default; host only)
as_anyrank(data, shape, stride, ndim, copy_meta) anyrank (inline store) copies into a TNY_MAX_RANK store (device-passable); as_anyrank<N>(…,copy_meta) sets capacity
as_anyrank<Space>(…) space-tagged anyrank Space = the data's memory space (default storage::view = host); pass storage::gpu_view for a device pointer so fixed/peel_front yield gpu_view views
as_anyrank(…, anyshape<etc,-1,-1,3>{}[, layout]) anyrank with a static tail bake the static TRAILING extents into the type (anyshape<etc,...>: etc = the erased batch, the dims after it = the static tail) — fixed/peel_front cells fold those inner dims (no per-call recast). An optional layout tag (ccontiguous/fcontiguous/strides<…>, like recast's 2nd arg; default keep_strides) also folds the trailing STRIDES — a contiguous inner block's cell mapping is then empty (EBO). Extents and strides are checked vs the runtime shape/strides once at the boundary, then trusted. Also on the copy_meta form. Bare anyshape<etc> ⇒ today's fully-dynamic carrier
as_anyrank(…, anyshape<3,etc,5>{}) anyrank with a static Head + Tail dims before etc are a static leading Head (anchored at 0), after it the Tailanyshape<3,etc,5> = (C_in=3, *spatial, C_out=5). The Head folds in fixed/dispatch_rank (full-rank window); peel_front<-Sr> stays trailing (the Head is peeled into the batch). Head extents fold; head strides stay runtime (a leading stride spans the dynamic middle). Empty Head (etc first) ⇒ the trailing-only carrier
from_dlpack<T[, Space]>(m) / from_dlpack<T, R[, Space]>(m) anyrank / rank-R view import; m may be a DLManagedTensor*, a bare DLTensor* (unmanaged — nothing to free, caller owns all lifetime), or a DLManagedTensorVersioned* (DLPack 1.0+). Space (default host) is checked against m→device — a kDLCUDA tensor needs storage::gpu_view
from_dlpack<T, anyshape<etc,-1,-1,3>[, Space]>(m[, layout]) anyrank with a static tail import with the static TRAILING extents baked in — the payload's trailing dims are checked against the tag at the import boundary, then every peeled cell folds them. Pass a layout tag by value (m, ccontiguous{}) to also fold the inner strides (checked vs the payload's strides — the "input is contiguous" precondition). Accepts all three carriers
to_dltensor(t, shape_out, strides_out[, dev]) DLTensor export to a bare DLTensor (no capsule/deleter/alloc): borrows t's data AND the two caller-supplied int64_t[≥rank] buffers it fills. Caller keeps the tensor + buffers alive. (to_dlpack(t) is the managed-capsule export.)
at.fixed<R>() rank-R dynamic_strides view requires ndim == R
dispatch_rank(at, f) bool call f with a fixed-rank view chosen by runtime ndim (one instantiation per total rank)
dispatch_rank<narrow_index>(at, f) bool ...and narrow each fixed cell's offset width to int32 when index_fits (rank outer, width inner; the leaf doubles). Default Narrow=false is the plain dispatch
dispatch_index<Idx2=int32_t>(v, f) void narrow a fixed-rank view v to Idx2 offsets when its span fits, else keep it — then call f; instantiates f for both widths
dispatch_layout(v, f) void the LAYOUT sibling: runtime-classify a dynamic_strides view as ccontiguous/fcontiguous/(else)dynamic_strides (is_dense<…>(), a stride compare) and call f with the retyped view — so recast<shape<-1,c,c>>() folds the inner strides safely. The C-order arm is the payoff (inner strides fold); f runs on up to 3 view types. Opt-in (triples instantiations)
at.peel_front<N>() / at.peel_front_at<N>(i) range / view batch idiom: keep the last \|N\| dims static, peel the rest (one kernel per \|N\|)
at.peel_front<-Sr>().enumerate() a range of {index, cell} range-for that ALSO yields the batch multi-index: for (auto [m, cell] : …) with m[d] = coord of batch axis d, m.rank() (runtime) = #batch axes, m.linear() = the flat batch index — for a per-batch-axis table param[d][m[d]]. Opt-in (the bare cell stays lean); m is a view of the live odometer (valid within the loop body). Raw iterator: it.index(d) / it.nbatch() / it.linear(). Composes with subrange
at.size_front<N>() → offset flattened batch count peel_front<N> yields (product of the peeled leading extents), no range built; N < 0
dispatch_value<Vs...>(v, f) bool call f(Int<k>{}) for the matching candidate k == v

Input → output type — fixed-rank views. The rank-erased carrier hands out a dyn_tensor<T, offset_t, R, Space> = tensor<T, cs::dextents<offset_t, R>, cs::layout_stride, Space> — a fully dynamic, arbitrarily-strided view whose Space is the carrier's view space (storage::view by default; storage::gpu_view for a device carrier):

Call Output type Notes
at.fixed<R>() dyn_tensor<T, offset_t, R, Space> requires ndim == R; extents + strides all runtime
at.peel_front_at<-Sr>(i) dyn_tensor<T, offset_t, Sr, Space> one batch cell (rank Sr); .recast<shape<-1,c,c>>() recovers static inner dims
at.peel_front_at<NewE[, NewL]>(i) / at.peel_front_at(i, shape<-1,c,c>{}[, ccontiguous{}]) cell with shape NewE fused peel+recast: the cell has the target trailing shape (rank = NewE::rank()) DIRECTLY — static inner extents fold, -1 stays dynamic, no separate recast. Removes the hand-kept Sr ≡ shape-rank invariant. Strides: keep_strides (default) keeps the carrier's runtime strides; a layout arg folds them (debug-checked promise), or dispatch_layout proves it
from_dlpack<T, R[, Space]>(m) dyn_tensor<T, offset_t, R, Space> rank-R view; Space checked against m→device
from_dlpack<T[, Space]>(m) anyrank (space-tagged) rank-erased; call fixed/peel_front/dispatch_rank on it

Recover static extents with .recast<NewE>() at the kernel boundary (see the recast type mapping) to fold the inner shape back into the type.

DLPack strides are in elements; numpy __array_interface__ in bytes (divide by the itemsize first).


Compile flags

Flag Effect
-DTNY_STD_PROMOTION standard C++ float promotion (wider wins) instead of lower-wins
-DTNY_NO_NEGATIVE_INDEX drop python-style negative-index wrap from operator() (tightest codegen)
-DTNY_PORTABLE_HALF force the portable software half/bfloat16 even under nvcc
-DTNY_MAX_RANK=N inline anyrank capacity / dispatch bound (default 32)
-DNDEBUG strip debug shape/precondition checks (host-only; already off on device)