Module Tolk.Renderer

Renderers: what code generation may emit for a target, and how it becomes a program.

A renderer describes one target to code generation: the launch sizes and memory it offers, the operations and data types it has natively, its tensor cores, the rewrites it needs last, how a linearized kernel is written as source, and the compiler that turns that source into a binary. Each target builds its renderer with v.

A renderer holds only facts its target fixes, never those of a device's physical part, such as its compute units or cache bytes: every device of a target runs one program of a kernel, which computes the same bits on each. A program is kept under its renderer's name and target (Codegen.program_key), so renderers of one name and target must make the same programs, and a renderer made with other facts takes another name.

Storage

#val with_storage : Ops.t -> Dtype.t -> Ops.t

with_storage u dt is the access u restated at dt on the storage it reaches through first sources: the Op.t.Param, Op.t.Buffer or Op.t.Alloc found from u by following first sources gets element type dt, and the nodes on the way are rebuilt on it. An access takes its data type from its storage, so accessing memory at another type retypes the storage.

Raises Invalid_argument if the first sources end before reaching a storage.

Cost estimates

#module Estimates : sig ... end

Cost estimates of kernels.

Compilers

#module Compiler : sig ... end

Compilers: source to binary.

Renderers

#type registers = {

#count : int;
(*
The number of vector registers.
*)
#bytes : int;
(*
The bytes of a vector register.
*)

}

The type for a target's vector register file.

#type t = private {

#name : string;
(*
The renderer's name, which tells apart renderers of one target.
*)
#target : Helpers.Target.t;
(*
The target rendered for.
*)
#suffix : string;
(*
A mark of the renderer's output, which tells apart programs of one device that different renderers made.
*)
#supports_float4 : bool;
(*
Loads and stores of consecutive elements may be merged into one vector access.
*)
#half8 : bool;
(*
Vectors of eight 16-bit floats are written, so that a copy of 16 bytes of them from global into local memory merges into one (Coalesce.memory_coalescing).
*)
#registers : unit -> registers option;
(*
registers () is the target's vector register file, if known, asked when first needed. It sets the widest merged access (Coalesce.memory_coalescing) and, on the CPU device, how many lanes a kernel's accumulators hold (Heuristic).
*)
#vector_lanes : int;
(*
The most lanes of a vector that arithmetic, logic, comparisons, selections and conversions compute at once: an operation of more lanes computes them in vectors of vector_lanes and a last of the rest, its last vector_lanes + 1 lanes as two vectors of vector_lanes / 2 + 1 and vector_lanes - vector_lanes / 2 where one lane would be left. At 1 or less, each lane is computed apart.
*)
#has_local : bool;
(*
The target launches workgroups of several threads, numbered by local indices.
*)
#has_shared : bool;
(*
The threads of a workgroup share memory.
*)
#global_max : int list;
(*
The greatest number of workgroups on each axis, x, y and z.
*)
#local_max : int list;
(*
The greatest number of threads of a workgroup on each axis.
*)
#global_prod_max : int list option;
(*
If any, the greatest number of threads on each axis, workgroups times their threads.
*)
#shared_max : int;
(*
The bytes of memory a workgroup shares.
*)
#async_copy : bool;
(*
A thread copies from global to shared memory without waiting for the copy, which a barrier then waits for: PTX's cp.async, from sm_80. Staged operands then take two stages (Stage.depth).
*)
#tensor_cores : Tc.t list;
(*
The tensor cores, in order of preference.
*)
#per_byte_first : int;
(*
The multiply-adds per byte of a contraction's streamed operand up to which it streams before a tensor core is tried (Heuristic): where a core's tiles start to run faster than threads that stream the operand, a crossover of the target's achieved rates. It moves the order of a sum, so it is fixed per target, never read from a device.
*)
#core_threads : int;
(*
The threads a kernel keeps once a tensor core's splits apply: a split applies only while core_threads / t workgroups of the core's t threads remain (Heuristic). Fewer leave memory short of loads in flight; more refuse reuse. It is a bound fitted to the products timed on one part of a target: 16384 on an RTX 5000 Ada (sm_89), and Metal's 32768 on an M1 Max, other Apple parts unmeasured. It moves the order of a sum, so it is fixed per target, never read from a device.
*)
#stream_threads : int;
(*
The threads that keep a part's memory busy, about as many loads in flight as its latency times its bandwidth. A kernel that streams an operand runs about as many (Heuristic). It is a bound fitted to the products timed on one part of a target: 32768 on an RTX 5000 Ada (sm_89), and Metal's 16384 on an M1 Max, other Apple parts unmeasured. Fixed per target, as core_threads is.
*)
#local_workgroups : int;
(*
The fewest workgroups a split of threads into larger workgroups leaves, so that each compute unit of a part takes about as many (Heuristic). A local split of a tensor core's N keeps the threads and packs a workgroup's warps onto one compute unit, which balances the units coarser: on an RTX 5000 Ada's 100 SMs a split that left 512 workgroups ran as fast as one that left 256, and 128 or 144 lost up to 1.34 times. It governs a core whose operands are not staged; a staged core's split shares each copied tile among its warps. Fixed per target, as core_threads is.
*)
#extra_matcher : (unit, Ops.t) Ops.Pattern_matcher.t;
(*
The rewrites the target needs after every other.
*)
#launch_members : int;
(*
The most members a launch's program holds (Shape.launch): a launch calls each member on its own, and the target's compiler takes longer for each past this count. A larger launch runs as several.
*)
#launch_args : int;
(*
The most arguments a launch passes: on CUDA, the 4 KiB of parameters a kernel takes, 512 of 8 bytes. A launch of more runs as several.
*)
#launch_locals : bool;
(*
Whether a member of a launch can declare workgroup memory. A member that fixes its threads or reduces may hold values there; where this is false, as on Metal, such a member runs alone.
*)
#packed_args : bool;
(*
A kernel takes its arguments as bytes: one parameter, which holds each argument at its offset (Device.Tiny_elf.iter_sig), a buffer by its address. The kernel loads each argument as its type, so the bytes must start on 8 bytes, the largest an argument takes.
*)
#code_for_op : (Op.t * (string list -> Dtype.t -> string)) list;
(*
The operations the target has natively, each with how it is written: f operands dtype is the expression of the operation on the expressions operands, producing dtype. Code generation decomposes the others.
*)
#native : Dtype.t -> bool;
(*
native dt is true iff the target has values of dt; see supported_dtypes.
*)
#fma : Dtype.t -> bool;
(*
fma dt holds only where an Op.t.Mulacc of the float dt, as the target writes it, rounds once to dt, subnormals aside on a target that flushes them. Code generation rewrites a float narrower than a float32 otherwise, unless its multiply-add in a float32 already rounds once (Codegen.full_rewrite_to_sink).
*)
#render : Ops.t list -> string;
(*
render uops is the source of the linearized kernel uops.
*)
#compiler : Compiler.t;
(*
The compiler of the rendered source.
*)

}

The type for renderers.

val v : 
  ?name:string ->
  ?suffix:string ->
  ?supports_float4:bool ->
  ?half8:bool ->
  ?registers:(unit -> registers option) ->
  ?vector_lanes:int ->
  ?has_local:bool ->
  ?has_shared:bool ->
  ?global_max:int list ->
  ?local_max:int list ->
  ?global_prod_max:int list ->
  ?shared_max:int ->
  ?async_copy:bool ->
  ?tensor_cores:Tc.t list ->
  ?per_byte_first:int ->
  ?core_threads:int ->
  ?stream_threads:int ->
  ?local_workgroups:int ->
  ?extra_matcher:(unit, Ops.t) Ops.Pattern_matcher.t ->
  ?launch_members:int ->
  ?launch_args:int ->
  ?launch_locals:bool ->
  ?packed_args:bool ->
  ?code_for_op:(Op.t * (string list -> Dtype.t -> string)) list ->
  ?native:(Dtype.t -> bool) ->
  ?fma:(Dtype.t -> bool) ->
  ?render:(Ops.t list -> string) ->
  ?compiler:Compiler.t ->
  Helpers.Target.t ->
  t

v target is the renderer for target with these fields. The defaults describe a target that renders nothing:

name is "Renderer" and suffix is "";
supports_float4, has_local and has_shared are true, registers gives None, half8 is false and vector_lanes is 1;
global_max and local_max are 0x8FFFFFFF on each axis, the greatest size a 32-bit signed index holds, and global_prod_max is None;
shared_max is 32768, and async_copy is false;
tensor_cores and code_for_op are empty, and so are the rules of extra_matcher;
per_byte_first is 64, sm_89's crossover, core_threads is 16384, stream_threads 32768 and local_workgroups 256, sm_89's, which every target takes until its own are measured;
launch_members is max_int, until a target's compiler is measured;
launch_args is max_int;
launch_locals is true;
packed_args is false;
native holds for every data type, and fma for Dtype.t.Float32 and Dtype.t.Float64;
render raises Invalid_argument;
compiler returns its source unchanged, and caches nothing.

#val with_compiler : Compiler.t -> t -> t

with_compiler c r is r with its rendered source compiled by c.

#val supported_dtypes : t -> Dtype.t list

supported_dtypes r is the data types of Dtype.all, in order, that r has natively, without Dtype.t.Float64 if Dtype.t.Int64 is emulated (setting Setting.emulated_dtypes): a double cannot be bitcast without a 64-bit integer.

Raises Invalid_argument if Setting.emulated_dtypes names no data type (Dtype.of_string).