# Runtime Dispatch

> NumKong for C

### Capabilities and Devices

Linking the library is the default, and the way to ship one binary across many CPU generations.
Every query writes its answer through a pointer and returns an `nk_status_t`.
CPU capabilities are reported along two independent axes, plus the mask to dispatch with:

- `nk_cpu_capabilities_detected` is what this CPU can execute, from CPUID or the OS, or where the OS cannot be asked, from what the compiler flags guarantee.
- `nk_cpu_capabilities_compiled` is what this binary contains, from the ISA probes at build time.
- `nk_cpu_capabilities_enabled` is both axes at once, the mask to pass every CPU call, and always holds `nk_cap_serial_k`.

Ask for `enabled` unless you specifically mean one of the raw axes.
The two axes are independent, and conflating them fails quietly rather than loudly: a binary whose ISA probes failed still reports this machine's full `detected` mask while containing no SIMD kernels at all.
To dispatch over fewer capabilities, pass a narrower mask to the call itself; there is no process-wide setting to change.

`nk_cpu_configure_thread` prepares the calling thread for the capabilities it is given, and only those.
Most capabilities need nothing: AMX on Linux costs one `arch_prctl` syscall, which grants tile state to the whole process, and fused BF16 dot products on Arm cost one `FPCR` write per thread.

```c
nk_capability_t enabled = nk_cap_serial_k;
nk_cpu_capabilities_enabled(&enabled);
nk_cpu_configure_thread(enabled);
if (enabled & nk_cap_sapphireamx_k) { /* AMX both detected and compiled in */ }
nk_dot_bf16_best(a, b, n, &dot, enabled & ~nk_cap_sapphireamx_k, NULL); // dispatch without AMX
```

GPUs are more ISAs with a mask of their own, one device at a time.
Each vendor has its own queries, and a device index is that runtime's own ordinal: the one `cudaSetDevice` or `hipSetDevice` takes, or the position in Metal's device list.
`nk_cuda_count_devices` counts the devices, and `nk_cuda_capabilities_detected`, `nk_cuda_capabilities_compiled`, and `nk_cuda_capabilities_enabled` report one device's capabilities the way their CPU twins do; `nk_rocm_*` and `nk_metal_*` are the same for the other vendors.
Each vendor's baseline, `nk_cap_cuda_k`, `nk_cap_rocm_k`, or `nk_cap_metal_k`, plays the role of `nk_cap_serial_k`.
A GPU capability, like `nk_dots_packed_bf16_ampere`, takes its CPU twin's arguments, queues on the trailing stream — a `cudaStream_t`, a `hipStream_t`, or an `id<MTLCommandQueue>`, with null for the default one — and returns without waiting.
Only the capability queries take an ordinal; the stream names the device for everything else.
`nk_memory_allocate_unified_best` and `nk_memory_free_unified_best` hand out and return memory the host and that device both address, and `nk_stream_synchronize_best` waits for the stream.
Batched kernels that may need scratch, like the packed and symmetric families and attention, take a `workspace, workspace_bytes` pair before the capabilities and never allocate.
Their `*_workspace_size_*` twin reports the bytes from the shapes alone, and a zero-byte workspace runs the scratch-free path where the capability has one.

```c
nk_size_t devices = 0, workspace_bytes = 0;
nk_capability_t gpu = 0;
if (nk_cuda_count_devices(&devices) == nk_success_k && nk_cuda_capabilities_enabled(0, &gpu) == nk_success_k) {
    nk_dots_packed_workspace_size_f64_best(rows, columns, depth, gpu, &workspace_bytes);
    void *workspace = NULL;
    nk_memory_allocate_unified_best(workspace_bytes, &workspace, gpu, cuda_stream);
    nk_dots_packed_f64_best(a, b_packed, c, rows, columns, depth, a_stride, c_stride, workspace, workspace_bytes, gpu,
                            cuda_stream);
    nk_stream_synchronize_best(gpu, cuda_stream);
}
```

`nk_capabilities_name` spells any such mask as the names bindings accept, like "serial,neon,neonhalf", into a buffer of `NUMKONG_CAPABILITIES_NAME_CAPACITY` bytes.

For exact register-level details, see `capabilities.h`.
Capability kernels, like `nk_dot_f32_haswell`, stay callable directly if you want to pin a path for testing or benchmarking.

`#include <numkong/numkong.h>`

`#include <numkong/capabilities.h>`

`#include <numkong/types.h>`

## nk_capability_t

64-bit bitmask representing SIMD capabilities of the target architecture.

```c
typedef nk_u64_t nk_capability_t;
```

## nk_kernel_kind_t

Enumeration of supported metric kinds.

```c
enum nk_kernel_kind_t
```

Some have aliases for convenience/discoverability.

- `nk_kernel_unknown_k` = `0`: Unknown kernel kind.
- `nk_kernel_dot_k` = `'i'`: Inner product.
- `nk_kernel_vdot_k` = `'v'`: Complex inner product.
- `nk_kernel_angular_k` = `'a'`: Angular (cosine) distance.
- `nk_kernel_euclidean_k` = `'e'`: Euclidean distance.
- `nk_kernel_sqeuclidean_k` = `'2'`: Squared Euclidean distance.
- `nk_kernel_hamming_k` = `'h'`: Hamming (or Manhattan) distance.
- `nk_kernel_jaccard_k` = `'j'`: Jaccard (or Tanimoto) coefficient.
- `nk_kernel_bilinear_k` = `'b'`: Bilinear form.
- `nk_kernel_mahalanobis_k` = `'m'`: Mahalanobis distance.
- `nk_kernel_kld_k` = `'k'`: Kullback-Leibler divergence.
- `nk_kernel_jsd_k` = `'s'`: Jensen-Shannon divergence.
- `nk_kernel_sparse_dot_k` = `'d'`: Sparse dot product with weighted indices.
- `nk_kernel_sparse_intersect_k` = `'x'`: Equivalent to unnormalized Jaccard.
- `nk_kernel_each_scale_k` = `'*'`: Element-wise Scale.
- `nk_kernel_each_sum_k` = `'+'`: Element-wise Sum.
- `nk_kernel_each_blend_k` = `'w'`: Element-wise Weighted Sum.
- `nk_kernel_each_fma_k` = `'f'`: Element-wise Fused Multiply-Add.
- `nk_kernel_each_swiglu_k` = `'W'`: Fused SwiGLU: y = silu(gate) ⊙ up (up = NULL → SiLU).
- `nk_kernel_each_rmsnorm_k` = `'H'`: Grouped RMSNorm: y = x · rsqrt(mean(x²) + epsilon) · γ.
- `nk_kernel_each_rmscast_k` = `'u'`: Grouped RMSNorm of F32 dot products into the narrower type they came from.
- `nk_kernel_trig_sin_k` = `'S'`: Element-wise sine.
- `nk_kernel_trig_cos_k` = `'C'`: Element-wise cosine.
- `nk_kernel_trig_atan_k` = `'A'`: Element-wise arctangent.
- `nk_kernel_reduce_moments_k` = `'R'`: Horizontal moments reduction (sum + sum-of-squares).
- `nk_kernel_reduce_minmax_k` = `'X'`: Horizontal minmax reduction (min + argmin + max + argmax).
- `nk_kernel_dots_pack_size_k` = `'P'`: GEMM packed buffer size.
- `nk_kernel_dots_pack_k` = `'Q'`: GEMM B matrix packing.
- `nk_kernel_dots_packed_k` = `'G'`: GEMM computation.
- `nk_kernel_dots_packed_shape_k` = `'J'`: GEMM packed shape read (width, depth).
- `nk_kernel_dots_symmetric_k` = `'y'`: Symmetric Gram matrix A × Aᵀ.
- `nk_kernel_dots_packed_workspace_size_k` = `'!'`: Workspace bytes of the GEMM computation.
- `nk_kernel_dots_symmetric_workspace_size_k` = `'#'`: Workspace bytes of the symmetric Gram matrix.
- `nk_kernel_angulars_packed_workspace_size_k` = `'$'`: Workspace bytes of batched angular distances (packed B).
- `nk_kernel_angulars_symmetric_workspace_size_k` = `'%'`: Workspace bytes of the symmetric angular distance matrix.
- `nk_kernel_euclideans_packed_workspace_size_k` = `'&'`: Workspace bytes of batched euclidean distances (packed B).
- `nk_kernel_euclideans_symmetric_workspace_size_k` = `','`: Workspace bytes of the symmetric euclidean distance matrix.
- `nk_kernel_hammings_packed_k` = `'M'`: Hamming distance computation.
- `nk_kernel_hammings_symmetric_k` = `'Y'`: Symmetric Hamming distance matrix A × Aᵀ.
- `nk_kernel_jaccards_packed_k` = `'p'`: Jaccard distance computation.
- `nk_kernel_jaccards_symmetric_k` = `'Z'`: Symmetric Jaccard distance matrix.
- `nk_kernel_hammings_packed_workspace_size_k` = `'.'`: Workspace bytes of the Hamming distance computation.
- `nk_kernel_hammings_symmetric_workspace_size_k` = `':'`: Workspace bytes of the symmetric Hamming distance matrix.
- `nk_kernel_jaccards_packed_workspace_size_k` = `';'`: Workspace bytes of the Jaccard distance computation.
- `nk_kernel_jaccards_symmetric_workspace_size_k` = `'='`: Workspace bytes of the symmetric Jaccard distance matrix.
- `nk_kernel_angulars_packed_k` = `'N'`: Batched angular distances (packed B).
- `nk_kernel_angulars_symmetric_k` = `'n'`: Symmetric angular distance matrix.
- `nk_kernel_euclideans_packed_k` = `'E'`: Batched euclidean distances (packed B).
- `nk_kernel_euclideans_symmetric_k` = `'D'`: Symmetric euclidean distance matrix.
- `nk_kernel_maxsim_pack_size_k` = `'L'`: MaxSim packed buffer size.
- `nk_kernel_maxsim_pack_k` = `'l'`: MaxSim vector packing.
- `nk_kernel_maxsim_packed_k` = `'T'`: MaxSim late-interaction computation.
- `nk_kernel_maxsim_packed_shape_k` = `'q'`: MaxSim packed shape read (vectors, depth).
- `nk_kernel_shapes_pack_size_k` = `'3'`: Packed cloud buffer size in bytes.
- `nk_kernel_shapes_pack_k` = `'4'`: Cloud packing: centering, principal axes and the pose header.
- `nk_kernel_shapes_pack_normalized_k` = `'7'`: Cloud packing that also divides each cloud by its RMS distance from the centroid.
- `nk_kernel_shapes_packed_shape_k` = `'0'`: Packed cloud read (points, centroid, frame, scale, spreads, largest radius).
- `nk_kernel_shapes_unpack_k` = `'1'`: Packed cloud restoration into the input type, the inverse of packing.
- `nk_kernel_chamfers_workspace_size_k` = `'~'`: Workspace bytes of many-to-many Chamfer.
- `nk_kernel_chamfers_packed_k` = `'8'`: Many-to-many Chamfer over principal-frame flips.
- `nk_kernel_hausdorffs_packed_k` = `'9'`: Many-to-many Hausdorff over principal-frame flips.
- `nk_kernel_wassersteinprocrustes_workspace_size_k` = `'5'`: Workspace bytes of many-to-many Wasserstein-Procrustes.
- `nk_kernel_wassersteinprocrustes_packed_k` = `'t'`: Many-to-many Wasserstein-Procrustes: exact transport alternated with Kabsch.
- `nk_kernel_gromovwassersteins_workspace_size_k` = `'6'`: Workspace bytes of many-to-many Gromov-Wasserstein.
- `nk_kernel_gromovwassersteins_packed_k` = `'c'`: Many-to-many Gromov-Wasserstein over plain distances within each cloud.
- `nk_kernel_rmsds_packed_k` = `'@'`: Many-to-many superposed RMSD between frames matched point by point.
- `nk_kernel_attention_pack_size_k` = `'B'`: Attention packed KV-cache Buffer size in Bytes.
- `nk_kernel_attention_pack_k` = `'V'`: Attention K/V packing (KV cache → the capability's layout).
- `nk_kernel_attention_packed_k` = `'F'`: Fused, Flash-style, attention under a band of visible keys.
- `nk_kernel_attention_packed_gradients_k` = `'g'`: Gradients of banded attention.
- `nk_kernel_attention_packed_workspace_size_k` = `'['`: Workspace bytes of banded attention.
- `nk_kernel_attention_packed_gradients_workspace_size_k` = `']'`: Workspace bytes of the gradients of banded attention.
- `nk_kernel_attention_packed_shape_k` = `'z'`: Attention packed shape read (heads, depth, segments).
- `nk_kernel_attention_rope_k` = `'I'`: NeoX split-half rotary position embedding (in place).
- `nk_kernel_sqeuclideans3d_k` = `'r'`: Many-to-many squared Euclidean distances between 3D points.
- `nk_kernel_haversines_pack_size_k` = `'('`: Haversine pack buffer size in bytes.
- `nk_kernel_haversines_pack_k` = `'<'`: Haversine packing: latitudes and longitudes into unit vectors.
- `nk_kernel_haversines_packed_k` = `'o'`: Many-to-many great-circle distances on the mean sphere.
- `nk_kernel_vincentys_pack_size_k` = `')'`: Vincenty pack buffer size in bytes.
- `nk_kernel_vincentys_pack_k` = `'>'`: Vincenty packing: longitudes and reduced latitudes.
- `nk_kernel_vincentys_packed_k` = `'O'`: Many-to-many geodesic distances on the ellipsoid.
- `nk_kernel_terrain_pack_size_k` = `'U'`: Terrain pack buffer size in bytes.
- `nk_kernel_terrain_pack_k` = `'K'`: Terrain packing: the maxima pyramid of a height grid.
- `nk_kernel_terrain_heights_packed_k` = `'^'`: Bilinear terrain heights under points.
- `nk_kernel_clearances_packed_k` = `'/'`: Many-to-many terrain clearances of sight lines.
- `nk_kernel_cast_k` = `'-'`: Type casting from one type to another.

## nk_kernel_punned_t

Any kernel, cast back to its kind's signature before the call.

```c
typedef void(*) nk_kernel_punned_t(void);
```

## nk_stream_t

A vendor stream: a `cudaStream_t`, `hipStream_t` or Metal command queue; null on CPUs.

```c
typedef struct nk_stream_opaque_ * nk_stream_t;
```

## nk_kernel_output_dtype

Returns the output dtype for a given metric kind and input dtype.

```c
nk_dtype_t nk_kernel_output_dtype(
    nk_kernel_kind_t kind,
    nk_dtype_t input);
```

## nk_kernel_is_commutative

Whether swapping the two inputs of `kind` leaves its result unchanged.

```c
int nk_kernel_is_commutative(nk_kernel_kind_t kind);
```

## nk_kernel_name

Canonical name of a kernel kind - the spelling bindings parse and interchange formats carry; "unknown" for unrecognized values.

```c
char const * nk_kernel_name(nk_kernel_kind_t kind);
```

## nk_kernel_named

Inverse of [`nk_kernel_name`](#nk_kernel_name) over an explicit-length string; [`nk_kernel_unknown_k`](#nk_kernel_unknown_k) for unrecognized names.

```c
nk_kernel_kind_t nk_kernel_named(char const *name, nk_size_t length);
```

## nk_cpu_capabilities_detected

Writes the CPU capabilities this machine can execute.

```c
nk_status_t nk_cpu_capabilities_detected(nk_capability_t *capabilities);
```

## nk_cpu_capabilities_compiled

Writes the CPU capabilities this binary holds kernels for.

```c
nk_status_t nk_cpu_capabilities_compiled(nk_capability_t *capabilities);
```

## nk_cpu_capabilities_enabled

Writes the CPU capabilities dispatch runs within: detected and compiled together.

```c
nk_status_t nk_cpu_capabilities_enabled(nk_capability_t *capabilities);
```

## nk_cpu_configure_thread

Prepares the calling thread for the kernels of `capabilities`, and only those.

```c
nk_status_t nk_cpu_configure_thread(nk_capability_t capabilities);
```

Most capabilities need nothing. The ones that do, and what they cost:

- [`nk_cap_sapphireamx_k`](#nk_cap_sapphireamx_k) on Linux: one `arch_prctl` syscall for AMX tile state, which the kernel grants to the whole process, and which the library already makes as it loads. Windows and FreeBSD need no request.
- Arm with FEAT_EBF16: one write of the thread's `FPCR`, making BF16 dot products fused. Finding FEAT_EBF16 costs a `sysctl` on Apple systems, and one ID-register read on Linux and FreeBSD when `capabilities` holds more than NEON.

- `capabilities` (in): The capabilities to prepare for, like [`nk_cpu_capabilities_enabled`](#nk_cpu_capabilities_enabled).

## nk_capabilities_name

Writes `capabilities` as a comma-separated name list such as "serial,haswell,skylake".

```c
nk_size_t nk_capabilities_name(
    nk_capability_t capabilities,
    char *buffer,
    nk_size_t capacity);
```

- `capabilities` (in): The mask to name, like `nk_cpu_capabilities_detected()` reports.
- `buffer` (out): Destination, always null-terminated; the list is truncated to fit.
- `capacity` (in): Size of `buffer`, like [`NUMKONG_CAPABILITIES_NAME_CAPACITY`](#NUMKONG_CAPABILITIES_NAME_CAPACITY); a zero capacity writes nothing.

## nk_stream_synchronize_best

Waits for everything queued on `stream`.

```c
nk_status_t nk_stream_synchronize_best(
    nk_capability_t capabilities,
    nk_stream_t stream);
```

- `nk_stream_synchronize` backends: serial

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k)

## nk_cuda_count_devices

Counts the CUDA devices the process sees.

```c
nk_status_t nk_cuda_count_devices(nk_size_t *count);
```

- `count` (out): The number of devices, zero without one.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Without one.

## nk_cuda_capabilities_detected

Reports the capabilities CUDA device `ordinal` runs.

```c
nk_status_t nk_cuda_capabilities_detected(
    nk_size_t ordinal,
    nk_capability_t *capabilities);
```

- `ordinal` (in): The CUDA runtime's ordinal, like `cudaSetDevice` takes.
- `capabilities` (out): The CUDA baseline and what the device runs, zero on failure.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Past the last device.

## nk_cuda_capabilities_compiled

Reports the CUDA capabilities this binary holds kernels for.

```c
nk_status_t nk_cuda_capabilities_compiled(
    nk_capability_t *capabilities);
```

- `capabilities` (out): The CUDA capabilities compiled in, zero without any.

## nk_cuda_capabilities_enabled

Reports the capabilities CUDA device `ordinal` runs that this binary holds kernels for.

```c
nk_status_t nk_cuda_capabilities_enabled(
    nk_size_t ordinal,
    nk_capability_t *capabilities);
```

- `ordinal` (in): The CUDA runtime's ordinal, like `cudaSetDevice` takes.
- `capabilities` (out): The CUDA baseline and what the device runs, zero on failure.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Past the last device.

## nk_stream_init_cuda

Creates a stream on CUDA device `ordinal` with `cudaStreamCreate`, for any consumer to use.

```c
nk_status_t nk_stream_init_cuda(nk_size_t ordinal, nk_stream_t *stream);
```

## nk_stream_free_cuda

Destroys a stream of [nk_stream_init_cuda](#nk_stream_init_cuda) with `cudaStreamDestroy`, once its work ends.

```c
nk_status_t nk_stream_free_cuda(nk_stream_t stream);
```

## nk_rocm_count_devices

Counts the ROCm devices the process sees.

```c
nk_status_t nk_rocm_count_devices(nk_size_t *count);
```

- `count` (out): The number of devices, zero without one.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Without one.

## nk_rocm_capabilities_detected

Reports the capabilities ROCm device `ordinal` runs.

```c
nk_status_t nk_rocm_capabilities_detected(
    nk_size_t ordinal,
    nk_capability_t *capabilities);
```

- `ordinal` (in): The HIP runtime's ordinal, like `hipSetDevice` takes.
- `capabilities` (out): The ROCm baseline and what the device runs, zero on failure.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Past the last device.

## nk_rocm_capabilities_compiled

Reports the ROCm capabilities this binary holds kernels for.

```c
nk_status_t nk_rocm_capabilities_compiled(
    nk_capability_t *capabilities);
```

- `capabilities` (out): The ROCm capabilities compiled in, zero without any.

## nk_rocm_capabilities_enabled

Reports the capabilities ROCm device `ordinal` runs that this binary holds kernels for.

```c
nk_status_t nk_rocm_capabilities_enabled(
    nk_size_t ordinal,
    nk_capability_t *capabilities);
```

- `ordinal` (in): The HIP runtime's ordinal, like `hipSetDevice` takes.
- `capabilities` (out): The ROCm baseline and what the device runs, zero on failure.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Past the last device.

## nk_stream_init_rocm

Creates a stream on ROCm device `ordinal` with `hipStreamCreate`, for any consumer to use.

```c
nk_status_t nk_stream_init_rocm(nk_size_t ordinal, nk_stream_t *stream);
```

## nk_stream_free_rocm

Destroys a stream of [nk_stream_init_rocm](#nk_stream_init_rocm) with `hipStreamDestroy`, once its work ends.

```c
nk_status_t nk_stream_free_rocm(nk_stream_t stream);
```

## nk_metal_count_devices

Counts the Metal devices the system lists.

```c
nk_status_t nk_metal_count_devices(nk_size_t *count);
```

- `count` (out): The number of devices, zero without one.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Without one.

## nk_metal_capabilities_detected

Reports the capabilities Metal device `ordinal` runs.

```c
nk_status_t nk_metal_capabilities_detected(
    nk_size_t ordinal,
    nk_capability_t *capabilities);
```

- `ordinal` (in): The device's position in the order the system lists them.
- `capabilities` (out): The Metal baseline when the device supports it, zero otherwise.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Past the last device.

## nk_metal_capabilities_compiled

Reports the Metal capabilities this binary holds kernels for.

```c
nk_status_t nk_metal_capabilities_compiled(
    nk_capability_t *capabilities);
```

- `capabilities` (out): The Metal capabilities compiled in, zero without any.

## nk_metal_capabilities_enabled

Reports the capabilities Metal device `ordinal` runs that this binary holds kernels for.

```c
nk_status_t nk_metal_capabilities_enabled(
    nk_size_t ordinal,
    nk_capability_t *capabilities);
```

- `ordinal` (in): The device's position in the order the system lists them.
- `capabilities` (out): The Metal baseline when the device supports it, zero otherwise.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) Past the last device.

## nk_stream_init_metal

Opens an `id<MTLCommandQueue>` stream on Metal device `ordinal` with `newCommandQueue`.

```c
nk_status_t nk_stream_init_metal(
    nk_size_t ordinal,
    nk_stream_t *stream);
```

## nk_stream_free_metal

Waits for a queue of [nk_stream_init_metal](#nk_stream_init_metal), then releases it.

```c
nk_status_t nk_stream_free_metal(nk_stream_t stream);
```

## nk_find_kernel_punned

Finds the kernel of `kind` for `dtype` that the best of `capabilities` runs.

```c
nk_status_t nk_find_kernel_punned(
    nk_kernel_kind_t kind,
    nk_dtype_t dtype,
    nk_capability_t capabilities,
    nk_kernel_punned_t *kernel,
    nk_capability_t *capability);
```

- `kind` (in): The kernel kind, like [`nk_kernel_dot_k`](#nk_kernel_dot_k).
- `dtype` (in): The input type, like [`nk_f32_k`](https://www.unum.cloud/docs/numkong/c/scalar-types.md#nk_f32_k).
- `capabilities` (in): One device's mask, like [`nk_cpu_capabilities_enabled`](#nk_cpu_capabilities_enabled) reports.
- `kernel` (out): The kernel, to cast back to its kind's signature, or null.
- `capability` (out): The single capability the kernel belongs to, or zero.

Statuses: - [`nk_missing_kernel_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_kernel_k) No capability in the mask has it.

## nk_cap_serial_k

Serial (non-SIMD) fallback capability.

```c
#define nk_cap_serial_k ((nk_capability_t)1)
```

Always available.

## nk_cap_any_k

Mask representing any capability.

```c
#define nk_cap_any_k ((nk_capability_t)NUMKONG_U64_MAX)
```

## nk_cap_neon_k

Arm NEON, the AArch64 Advanced SIMD baseline.

```c
#define nk_cap_neon_k ((nk_capability_t)1 <<1)
```

## nk_cap_neonhalf_k

NEON with half-precision arithmetic, FEAT_FP16.

```c
#define nk_cap_neonhalf_k ((nk_capability_t)1 <<2)
```

## nk_cap_neonbfdot_k

NEON with BFloat16 dot products, FEAT_BF16.

```c
#define nk_cap_neonbfdot_k ((nk_capability_t)1 <<3)
```

## nk_cap_neonfhm_k

NEON with half-precision multiply-accumulate into single precision, FEAT_FHM.

```c
#define nk_cap_neonfhm_k ((nk_capability_t)1 <<4)
```

## nk_cap_neonsdot_k

NEON with 8-bit integer dot products, FEAT_DotProd.

```c
#define nk_cap_neonsdot_k ((nk_capability_t)1 <<5)
```

## nk_cap_neonfp8_k

NEON with FP8 dot products, FEAT_FP8DOT4.

```c
#define nk_cap_neonfp8_k ((nk_capability_t)1 <<6)
```

## nk_cap_sve_k

Arm Scalable Vector Extension.

```c
#define nk_cap_sve_k ((nk_capability_t)1 <<7)
```

## nk_cap_svehalf_k

SVE with half-precision arithmetic.

```c
#define nk_cap_svehalf_k ((nk_capability_t)1 <<8)
```

## nk_cap_svesdot_k

SVE with 8-bit integer matrix multiplies, FEAT_I8MM.

```c
#define nk_cap_svesdot_k ((nk_capability_t)1 <<9)
```

## nk_cap_svebfdot_k

SVE with BFloat16 dot products.

```c
#define nk_cap_svebfdot_k ((nk_capability_t)1 <<10)
```

## nk_cap_sve2_k

Arm SVE2.

```c
#define nk_cap_sve2_k ((nk_capability_t)1 <<11)
```

## nk_cap_sme_k

Arm Scalable Matrix Extension.

```c
#define nk_cap_sme_k ((nk_capability_t)1 <<12)
```

## nk_cap_smef64_k

SME with double-precision outer products, FEAT_SME_F64F64.

```c
#define nk_cap_smef64_k ((nk_capability_t)1 <<13)
```

## nk_cap_smebi32_k

SME with binary outer products into 32-bit integers, SME_BI32I32.

```c
#define nk_cap_smebi32_k ((nk_capability_t)1 <<14)
```

## nk_cap_haswell_k

x86 AVX2, FMA, F16C, BMI1 and BMI2, as on Intel Haswell.

```c
#define nk_cap_haswell_k ((nk_capability_t)1 <<15)
```

## nk_cap_alder_k

[`nk_cap_haswell_k`](#nk_cap_haswell_k) plus AVX-VNNI, as on Intel Alder Lake.

```c
#define nk_cap_alder_k ((nk_capability_t)1 <<16)
```

## nk_cap_sierra_k

[`nk_cap_alder_k`](#nk_cap_alder_k) plus AVX-VNNI-INT8, as on Intel Sierra Forest.

```c
#define nk_cap_sierra_k ((nk_capability_t)1 <<17)
```

## nk_cap_skylake_k

[`nk_cap_haswell_k`](#nk_cap_haswell_k) plus AVX-512 F, DQ, BW and VL, as on Intel Skylake-X.

```c
#define nk_cap_skylake_k ((nk_capability_t)1 <<18)
```

## nk_cap_icelake_k

[`nk_cap_skylake_k`](#nk_cap_skylake_k) plus AVX-512 VNNI, IFMA, BITALG, VBMI, VBMI2 and VPOPCNTDQ.

```c
#define nk_cap_icelake_k ((nk_capability_t)1 <<19)
```

## nk_cap_genoa_k

[`nk_cap_skylake_k`](#nk_cap_skylake_k) plus AVX-512 BF16 and VNNI, as on AMD Genoa.

```c
#define nk_cap_genoa_k ((nk_capability_t)1 <<20)
```

## nk_cap_turin_k

[`nk_cap_genoa_k`](#nk_cap_genoa_k) plus AVX-512 VBMI2 and VP2INTERSECT, as on AMD Turin.

```c
#define nk_cap_turin_k ((nk_capability_t)1 <<21)
```

## nk_cap_sapphire_k

[`nk_cap_skylake_k`](#nk_cap_skylake_k) plus AVX-512 FP16, as on Intel Sapphire Rapids.

```c
#define nk_cap_sapphire_k ((nk_capability_t)1 <<22)
```

## nk_cap_diamond_k

[`nk_cap_sapphire_k`](#nk_cap_sapphire_k) plus AVX10.2, as on Intel Diamond Rapids.

```c
#define nk_cap_diamond_k ((nk_capability_t)1 <<23)
```

## nk_cap_sapphireamx_k

[`nk_cap_sapphire_k`](#nk_cap_sapphire_k) and [`nk_cap_genoa_k`](#nk_cap_genoa_k) plus AMX tiles with BF16 and INT8.

```c
#define nk_cap_sapphireamx_k ((nk_capability_t)1 <<24)
```

## nk_cap_graniteamx_k

[`nk_cap_sapphireamx_k`](#nk_cap_sapphireamx_k) plus AMX-FP16, as on Intel Granite Rapids.

```c
#define nk_cap_graniteamx_k ((nk_capability_t)1 <<25)
```

## nk_cap_diamondamx_k

[`nk_cap_sapphireamx_k`](#nk_cap_sapphireamx_k) plus AMX-FP8 and AMX-AVX512, as on Intel Diamond Rapids.

```c
#define nk_cap_diamondamx_k ((nk_capability_t)1 <<26)
```

## nk_cap_rvv_k

RISC-V Vector extension.

```c
#define nk_cap_rvv_k ((nk_capability_t)1 <<27)
```

## nk_cap_rvvbf16_k

RVV with BFloat16 widening multiply-accumulate, Zvfbfwma.

```c
#define nk_cap_rvvbf16_k ((nk_capability_t)1 <<28)
```

## nk_cap_rvvhalf_k

RVV with half-precision arithmetic, Zvfh.

```c
#define nk_cap_rvvhalf_k ((nk_capability_t)1 <<29)
```

## nk_cap_rvvbb_k

RVV with basic bit manipulation, Zvbb.

```c
#define nk_cap_rvvbb_k ((nk_capability_t)1 <<30)
```

## nk_cap_v128_k

WebAssembly 128-bit SIMD.

```c
#define nk_cap_v128_k ((nk_capability_t)1 <<31)
```

## nk_cap_v128relaxed_k

WebAssembly Relaxed SIMD.

```c
#define nk_cap_v128relaxed_k ((nk_capability_t)1 <<32)
```

## nk_cap_powervsx_k

IBM Power VSX at ISA 3.0, as on POWER9.

```c
#define nk_cap_powervsx_k ((nk_capability_t)1 <<33)
```

## nk_cap_loongsonasx_k

LoongArch LASX 256-bit SIMD.

```c
#define nk_cap_loongsonasx_k ((nk_capability_t)1 <<34)
```

## nk_cap_cuda_k

NVIDIA baseline, which runs on every CUDA device.

```c
#define nk_cap_cuda_k ((nk_capability_t)1 <<48)
```

## nk_cap_ampere_k

NVIDIA `mma.sync` tensor cores, as on A100, RTX 30 and every later GPU.

```c
#define nk_cap_ampere_k ((nk_capability_t)1 <<49)
```

## nk_cap_ada_k

NVIDIA FP8 `mma.sync`, as on L4, L40S, RTX 40 and every later GPU.

```c
#define nk_cap_ada_k ((nk_capability_t)1 <<50)
```

## nk_cap_hopper_k

NVIDIA Hopper `wgmma`, as on H100, H200 and GH200.

```c
#define nk_cap_hopper_k ((nk_capability_t)1 <<51)
```

## nk_cap_blackwell_k

NVIDIA datacenter Blackwell `tcgen05`, as on B200, GB200, B300 and GB300.

```c
#define nk_cap_blackwell_k ((nk_capability_t)1 <<52)
```

## nk_cap_blackwellrtx_k

NVIDIA FP6 and FP4 `mma.sync` of compute capability 12.x, as on RTX 50 and DGX Spark.

```c
#define nk_cap_blackwellrtx_k ((nk_capability_t)1 <<53)
```

## nk_cap_blackwellultra_k

NVIDIA Blackwell Ultra `tcgen05.ld.red`, as on B300 and GB300.

```c
#define nk_cap_blackwellultra_k ((nk_capability_t)1 <<54)
```

## nk_cap_rocm_k

AMD baseline, which runs on every ROCm device.

```c
#define nk_cap_rocm_k ((nk_capability_t)1 <<56)
```

## nk_cap_cdna3_k

AMD MFMA over BF16, F16 and I8, as on MI300X, MI300A and MI325X.

```c
#define nk_cap_cdna3_k ((nk_capability_t)1 <<57)
```

## nk_cap_cdna4_k

AMD MFMA over FP8, FP6 and FP4, as on MI350X and MI355X.

```c
#define nk_cap_cdna4_k ((nk_capability_t)1 <<58)
```

## nk_cap_cdna5_k

AMD WMMA over FP8, FP6 and FP4, as on MI450X, MI455X and MI430X.

```c
#define nk_cap_cdna5_k ((nk_capability_t)1 <<59)
```

## nk_cap_metal_k

Apple baseline, which runs on every Metal device.

```c
#define nk_cap_metal_k ((nk_capability_t)1 <<60)
```

## nk_cap_apple9_k

Apple GPU family 9, as on M3 and M4.

```c
#define nk_cap_apple9_k ((nk_capability_t)1 <<61)
```

## nk_cap_apple10_k

Apple GPU family 10, as on M5.

```c
#define nk_cap_apple10_k ((nk_capability_t)1 <<62)
```

## nk_cap_gpus_k

Every GPU capability above, which the CPU dispatch never detects, compiles or enables.

```c
#define nk_cap_gpus_k (nk_cap_cuda_k | nk_cap_ampere_k | nk_cap_ada_k | nk_cap_hopper_k | nk_cap_blackwell_k | nk_cap_blackwellrtx_k | \ nk_cap_blackwellultra_k | nk_cap_rocm_k | nk_cap_cdna3_k | nk_cap_cdna4_k | nk_cap_cdna5_k | nk_cap_metal_k | \ nk_cap_apple9_k | nk_cap_apple10_k)
```

## nk_cap_cpus_k

Every CPU capability, the bits below the first GPU vendor's.

```c
#define nk_cap_cpus_k (nk_cap_cuda_k - 1)
```

## NUMKONG_CAPABILITIES_NAME_CAPACITY

Buffer size [`nk_capabilities_name`](#nk_capabilities_name) never overruns, including its null terminator.

```c
#define NUMKONG_CAPABILITIES_NAME_CAPACITY 1024
```