C · Runtime Dispatch

Runtime Dispatch

#include <numkong/numkong.h>#include <numkong/capabilities.h>#include <numkong/types.h> 8.0.0-rc.1 · b1f866ab

§ 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.

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.

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.

API

Types

nk_capability_t64-bit bitmask representing SIMD capabilities of the target architecture.
capabilities.h:188 ↗
typedef nk_u64_t nk_capability_t;
nk_kernel_kind_tEnumeration of supported metric kinds.
capabilities.h:433 ↗
enum nk_kernel_kind_t

Some have aliases for convenience/discoverability.

Values

nk_kernel_unknown_k0Unknown 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_tAny kernel, cast back to its kind's signature before the call.
capabilities.h:1102 ↗
typedef void(*) nk_kernel_punned_t(void);
nk_stream_tA vendor stream: a cudaStream_t, hipStream_t or Metal command queue; null on CPUs.
types.h:1445 ↗
typedef struct nk_stream_opaque_ * nk_stream_t;

Functions

nk_kernel_output_dtypeReturns the output dtype for a given metric kind and input dtype.
numkong.h:40 ↗
nk_dtype_t nk_kernel_output_dtype(
    nk_kernel_kind_t kind,
    nk_dtype_t input);
nk_kernel_is_commutativeWhether swapping the two inputs of kind leaves its result unchanged.
numkong.h:82 ↗
int nk_kernel_is_commutative(nk_kernel_kind_t kind);
nk_kernel_nameCanonical name of a kernel kind - the spelling bindings parse and interchange formats carry; "unknown" for unrecognized values.
capabilities.h:696 ↗
char const * nk_kernel_name(nk_kernel_kind_t kind);
nk_kernel_namedInverse of nk_kernel_name over an explicit-length string; nk_kernel_unknown_k for unrecognized names.
capabilities.h:790 ↗
nk_kernel_kind_t nk_kernel_named(char const *name, nk_size_t length);
nk_cpu_capabilities_detectedWrites the CPU capabilities this machine can execute.
capabilities.h:1648 ↗
nk_status_t nk_cpu_capabilities_detected(nk_capability_t *capabilities);
nk_cpu_capabilities_compiledWrites the CPU capabilities this binary holds kernels for.
capabilities.h:1651 ↗
nk_status_t nk_cpu_capabilities_compiled(nk_capability_t *capabilities);
nk_cpu_capabilities_enabledWrites the CPU capabilities dispatch runs within: detected and compiled together.
capabilities.h:1654 ↗
nk_status_t nk_cpu_capabilities_enabled(nk_capability_t *capabilities);
nk_cpu_configure_threadPrepares the calling thread for the kernels of capabilities, and only those.
capabilities.h:1668 ↗
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 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.

Parameters

capabilitiesinThe capabilities to prepare for, like nk_cpu_capabilities_enabled.
nk_capabilities_nameWrites capabilities as a comma-separated name list such as "serial,haswell,skylake".
capabilities.h:1678 ↗
nk_size_t nk_capabilities_name(
    nk_capability_t capabilities,
    char *buffer,
    nk_size_t capacity);

Parameters

capabilitiesinThe mask to name, like nk_cpu_capabilities_detected() reports.
bufferoutDestination, always null-terminated; the list is truncated to fit.
capacityinSize of buffer, like NUMKONG_CAPABILITIES_NAME_CAPACITY; a zero capacity writes nothing.

Returns

Bytes written, excluding the null terminator.

nk_stream_synchronize_bestWaits for everything queued on stream.
capabilities.h:1714 ↗
nk_status_t nk_stream_synchronize_best(
    nk_capability_t capabilities,
    nk_stream_t stream);

Backends

  • nk_stream_synchronizeserial

Statuses

nk_cuda_count_devicesCounts the CUDA devices the process sees.
capabilities.h:1832 ↗
nk_status_t nk_cuda_count_devices(nk_size_t *count);

Parameters

countoutThe number of devices, zero without one.

Statuses

nk_cuda_capabilities_detectedReports the capabilities CUDA device ordinal runs.
capabilities.h:1842 ↗
nk_status_t nk_cuda_capabilities_detected(
    nk_size_t ordinal,
    nk_capability_t *capabilities);

Parameters

ordinalinThe CUDA runtime's ordinal, like cudaSetDevice takes.
capabilitiesoutThe CUDA baseline and what the device runs, zero on failure.

Statuses

nk_cuda_capabilities_compiledReports the CUDA capabilities this binary holds kernels for.
capabilities.h:1848 ↗
nk_status_t nk_cuda_capabilities_compiled(
    nk_capability_t *capabilities);

Parameters

capabilitiesoutThe CUDA capabilities compiled in, zero without any.
nk_cuda_capabilities_enabledReports the capabilities CUDA device ordinal runs that this binary holds kernels for.
capabilities.h:1858 ↗
nk_status_t nk_cuda_capabilities_enabled(
    nk_size_t ordinal,
    nk_capability_t *capabilities);

Parameters

ordinalinThe CUDA runtime's ordinal, like cudaSetDevice takes.
capabilitiesoutThe CUDA baseline and what the device runs, zero on failure.

Statuses

nk_stream_init_cudaCreates a stream on CUDA device ordinal with cudaStreamCreate, for any consumer to use.
capabilities.h:1861 ↗
nk_status_t nk_stream_init_cuda(nk_size_t ordinal, nk_stream_t *stream);
nk_stream_free_cudaDestroys a stream of nk_stream_init_cuda with cudaStreamDestroy, once its work ends.
capabilities.h:1864 ↗
nk_status_t nk_stream_free_cuda(nk_stream_t stream);
nk_rocm_count_devicesCounts the ROCm devices the process sees.
capabilities.h:1873 ↗
nk_status_t nk_rocm_count_devices(nk_size_t *count);

Parameters

countoutThe number of devices, zero without one.

Statuses

nk_rocm_capabilities_detectedReports the capabilities ROCm device ordinal runs.
capabilities.h:1883 ↗
nk_status_t nk_rocm_capabilities_detected(
    nk_size_t ordinal,
    nk_capability_t *capabilities);

Parameters

ordinalinThe HIP runtime's ordinal, like hipSetDevice takes.
capabilitiesoutThe ROCm baseline and what the device runs, zero on failure.

Statuses

nk_rocm_capabilities_compiledReports the ROCm capabilities this binary holds kernels for.
capabilities.h:1889 ↗
nk_status_t nk_rocm_capabilities_compiled(
    nk_capability_t *capabilities);

Parameters

capabilitiesoutThe ROCm capabilities compiled in, zero without any.
nk_rocm_capabilities_enabledReports the capabilities ROCm device ordinal runs that this binary holds kernels for.
capabilities.h:1899 ↗
nk_status_t nk_rocm_capabilities_enabled(
    nk_size_t ordinal,
    nk_capability_t *capabilities);

Parameters

ordinalinThe HIP runtime's ordinal, like hipSetDevice takes.
capabilitiesoutThe ROCm baseline and what the device runs, zero on failure.

Statuses

nk_stream_init_rocmCreates a stream on ROCm device ordinal with hipStreamCreate, for any consumer to use.
capabilities.h:1902 ↗
nk_status_t nk_stream_init_rocm(nk_size_t ordinal, nk_stream_t *stream);
nk_stream_free_rocmDestroys a stream of nk_stream_init_rocm with hipStreamDestroy, once its work ends.
capabilities.h:1905 ↗
nk_status_t nk_stream_free_rocm(nk_stream_t stream);
nk_metal_count_devicesCounts the Metal devices the system lists.
capabilities.h:1914 ↗
nk_status_t nk_metal_count_devices(nk_size_t *count);

Parameters

countoutThe number of devices, zero without one.

Statuses

nk_metal_capabilities_detectedReports the capabilities Metal device ordinal runs.
capabilities.h:1924 ↗
nk_status_t nk_metal_capabilities_detected(
    nk_size_t ordinal,
    nk_capability_t *capabilities);

Parameters

ordinalinThe device's position in the order the system lists them.
capabilitiesoutThe Metal baseline when the device supports it, zero otherwise.

Statuses

nk_metal_capabilities_compiledReports the Metal capabilities this binary holds kernels for.
capabilities.h:1930 ↗
nk_status_t nk_metal_capabilities_compiled(
    nk_capability_t *capabilities);

Parameters

capabilitiesoutThe Metal capabilities compiled in, zero without any.
nk_metal_capabilities_enabledReports the capabilities Metal device ordinal runs that this binary holds kernels for.
capabilities.h:1940 ↗
nk_status_t nk_metal_capabilities_enabled(
    nk_size_t ordinal,
    nk_capability_t *capabilities);

Parameters

ordinalinThe device's position in the order the system lists them.
capabilitiesoutThe Metal baseline when the device supports it, zero otherwise.

Statuses

nk_stream_init_metalOpens an id<MTLCommandQueue> stream on Metal device ordinal with newCommandQueue.
capabilities.h:1943 ↗
nk_status_t nk_stream_init_metal(
    nk_size_t ordinal,
    nk_stream_t *stream);
nk_stream_free_metalWaits for a queue of nk_stream_init_metal, then releases it.
capabilities.h:1946 ↗
nk_status_t nk_stream_free_metal(nk_stream_t stream);
nk_find_kernel_punnedFinds the kernel of kind for dtype that the best of capabilities runs.
capabilities.h:1959 ↗
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);

Parameters

kindinThe kernel kind, like nk_kernel_dot_k.
dtypeinThe input type, like nk_f32_k.
capabilitiesinOne device's mask, like nk_cpu_capabilities_enabled reports.
kerneloutThe kernel, to cast back to its kind's signature, or null.
capabilityoutThe single capability the kernel belongs to, or zero.

Statuses

Macros

nk_cap_serial_kSerial (non-SIMD) fallback capability.
capabilities.h:191 ↗
#define nk_cap_serial_k ((nk_capability_t)1)

Always available.

nk_cap_any_kMask representing any capability.
capabilities.h:194 ↗
#define nk_cap_any_k ((nk_capability_t)NUMKONG_U64_MAX)
nk_cap_neon_kArm NEON, the AArch64 Advanced SIMD baseline.
capabilities.h:200 ↗
#define nk_cap_neon_k ((nk_capability_t)1 <<1)
nk_cap_neonhalf_kNEON with half-precision arithmetic, FEAT_FP16.
capabilities.h:203 ↗
#define nk_cap_neonhalf_k ((nk_capability_t)1 <<2)
nk_cap_neonbfdot_kNEON with BFloat16 dot products, FEAT_BF16.
capabilities.h:206 ↗
#define nk_cap_neonbfdot_k ((nk_capability_t)1 <<3)
nk_cap_neonfhm_kNEON with half-precision multiply-accumulate into single precision, FEAT_FHM.
capabilities.h:209 ↗
#define nk_cap_neonfhm_k ((nk_capability_t)1 <<4)
nk_cap_neonsdot_kNEON with 8-bit integer dot products, FEAT_DotProd.
capabilities.h:212 ↗
#define nk_cap_neonsdot_k ((nk_capability_t)1 <<5)
nk_cap_neonfp8_kNEON with FP8 dot products, FEAT_FP8DOT4.
capabilities.h:215 ↗
#define nk_cap_neonfp8_k ((nk_capability_t)1 <<6)
nk_cap_sve_kArm Scalable Vector Extension.
capabilities.h:218 ↗
#define nk_cap_sve_k ((nk_capability_t)1 <<7)
nk_cap_svehalf_kSVE with half-precision arithmetic.
capabilities.h:221 ↗
#define nk_cap_svehalf_k ((nk_capability_t)1 <<8)
nk_cap_svesdot_kSVE with 8-bit integer matrix multiplies, FEAT_I8MM.
capabilities.h:224 ↗
#define nk_cap_svesdot_k ((nk_capability_t)1 <<9)
nk_cap_svebfdot_kSVE with BFloat16 dot products.
capabilities.h:227 ↗
#define nk_cap_svebfdot_k ((nk_capability_t)1 <<10)
nk_cap_sve2_kArm SVE2.
capabilities.h:230 ↗
#define nk_cap_sve2_k ((nk_capability_t)1 <<11)
nk_cap_sme_kArm Scalable Matrix Extension.
capabilities.h:233 ↗
#define nk_cap_sme_k ((nk_capability_t)1 <<12)
nk_cap_smef64_kSME with double-precision outer products, FEAT_SME_F64F64.
capabilities.h:236 ↗
#define nk_cap_smef64_k ((nk_capability_t)1 <<13)
nk_cap_smebi32_kSME with binary outer products into 32-bit integers, SME_BI32I32.
capabilities.h:239 ↗
#define nk_cap_smebi32_k ((nk_capability_t)1 <<14)
nk_cap_haswell_kx86 AVX2, FMA, F16C, BMI1 and BMI2, as on Intel Haswell.
capabilities.h:242 ↗
#define nk_cap_haswell_k ((nk_capability_t)1 <<15)
nk_cap_alder_knk_cap_haswell_k plus AVX-VNNI, as on Intel Alder Lake.
capabilities.h:245 ↗
#define nk_cap_alder_k ((nk_capability_t)1 <<16)
nk_cap_sierra_knk_cap_alder_k plus AVX-VNNI-INT8, as on Intel Sierra Forest.
capabilities.h:248 ↗
#define nk_cap_sierra_k ((nk_capability_t)1 <<17)
nk_cap_skylake_knk_cap_haswell_k plus AVX-512 F, DQ, BW and VL, as on Intel Skylake-X.
capabilities.h:251 ↗
#define nk_cap_skylake_k ((nk_capability_t)1 <<18)
nk_cap_icelake_knk_cap_skylake_k plus AVX-512 VNNI, IFMA, BITALG, VBMI, VBMI2 and VPOPCNTDQ.
capabilities.h:254 ↗
#define nk_cap_icelake_k ((nk_capability_t)1 <<19)
nk_cap_genoa_knk_cap_skylake_k plus AVX-512 BF16 and VNNI, as on AMD Genoa.
capabilities.h:257 ↗
#define nk_cap_genoa_k ((nk_capability_t)1 <<20)
nk_cap_turin_knk_cap_genoa_k plus AVX-512 VBMI2 and VP2INTERSECT, as on AMD Turin.
capabilities.h:260 ↗
#define nk_cap_turin_k ((nk_capability_t)1 <<21)
nk_cap_sapphire_knk_cap_skylake_k plus AVX-512 FP16, as on Intel Sapphire Rapids.
capabilities.h:263 ↗
#define nk_cap_sapphire_k ((nk_capability_t)1 <<22)
nk_cap_diamond_knk_cap_sapphire_k plus AVX10.2, as on Intel Diamond Rapids.
capabilities.h:266 ↗
#define nk_cap_diamond_k ((nk_capability_t)1 <<23)
nk_cap_sapphireamx_knk_cap_sapphire_k and nk_cap_genoa_k plus AMX tiles with BF16 and INT8.
capabilities.h:269 ↗
#define nk_cap_sapphireamx_k ((nk_capability_t)1 <<24)
nk_cap_graniteamx_knk_cap_sapphireamx_k plus AMX-FP16, as on Intel Granite Rapids.
capabilities.h:272 ↗
#define nk_cap_graniteamx_k ((nk_capability_t)1 <<25)
nk_cap_diamondamx_knk_cap_sapphireamx_k plus AMX-FP8 and AMX-AVX512, as on Intel Diamond Rapids.
capabilities.h:275 ↗
#define nk_cap_diamondamx_k ((nk_capability_t)1 <<26)
nk_cap_rvv_kRISC-V Vector extension.
capabilities.h:278 ↗
#define nk_cap_rvv_k ((nk_capability_t)1 <<27)
nk_cap_rvvbf16_kRVV with BFloat16 widening multiply-accumulate, Zvfbfwma.
capabilities.h:281 ↗
#define nk_cap_rvvbf16_k ((nk_capability_t)1 <<28)
nk_cap_rvvhalf_kRVV with half-precision arithmetic, Zvfh.
capabilities.h:284 ↗
#define nk_cap_rvvhalf_k ((nk_capability_t)1 <<29)
nk_cap_rvvbb_kRVV with basic bit manipulation, Zvbb.
capabilities.h:287 ↗
#define nk_cap_rvvbb_k ((nk_capability_t)1 <<30)
nk_cap_v128_kWebAssembly 128-bit SIMD.
capabilities.h:290 ↗
#define nk_cap_v128_k ((nk_capability_t)1 <<31)
nk_cap_v128relaxed_kWebAssembly Relaxed SIMD.
capabilities.h:293 ↗
#define nk_cap_v128relaxed_k ((nk_capability_t)1 <<32)
nk_cap_powervsx_kIBM Power VSX at ISA 3.0, as on POWER9.
capabilities.h:296 ↗
#define nk_cap_powervsx_k ((nk_capability_t)1 <<33)
nk_cap_loongsonasx_kLoongArch LASX 256-bit SIMD.
capabilities.h:299 ↗
#define nk_cap_loongsonasx_k ((nk_capability_t)1 <<34)
nk_cap_cuda_kNVIDIA baseline, which runs on every CUDA device.
capabilities.h:306 ↗
#define nk_cap_cuda_k ((nk_capability_t)1 <<48)
nk_cap_ampere_kNVIDIA mma.sync tensor cores, as on A100, RTX 30 and every later GPU.
capabilities.h:309 ↗
#define nk_cap_ampere_k ((nk_capability_t)1 <<49)
nk_cap_ada_kNVIDIA FP8 mma.sync, as on L4, L40S, RTX 40 and every later GPU.
capabilities.h:312 ↗
#define nk_cap_ada_k ((nk_capability_t)1 <<50)
nk_cap_hopper_kNVIDIA Hopper wgmma, as on H100, H200 and GH200.
capabilities.h:315 ↗
#define nk_cap_hopper_k ((nk_capability_t)1 <<51)
nk_cap_blackwell_kNVIDIA datacenter Blackwell tcgen05, as on B200, GB200, B300 and GB300.
capabilities.h:318 ↗
#define nk_cap_blackwell_k ((nk_capability_t)1 <<52)
nk_cap_blackwellrtx_kNVIDIA FP6 and FP4 mma.sync of compute capability 12.x, as on RTX 50 and DGX Spark.
capabilities.h:321 ↗
#define nk_cap_blackwellrtx_k ((nk_capability_t)1 <<53)
nk_cap_blackwellultra_kNVIDIA Blackwell Ultra tcgen05.ld.red, as on B300 and GB300.
capabilities.h:324 ↗
#define nk_cap_blackwellultra_k ((nk_capability_t)1 <<54)
nk_cap_rocm_kAMD baseline, which runs on every ROCm device.
capabilities.h:327 ↗
#define nk_cap_rocm_k ((nk_capability_t)1 <<56)
nk_cap_cdna3_kAMD MFMA over BF16, F16 and I8, as on MI300X, MI300A and MI325X.
capabilities.h:330 ↗
#define nk_cap_cdna3_k ((nk_capability_t)1 <<57)
nk_cap_cdna4_kAMD MFMA over FP8, FP6 and FP4, as on MI350X and MI355X.
capabilities.h:333 ↗
#define nk_cap_cdna4_k ((nk_capability_t)1 <<58)
nk_cap_cdna5_kAMD WMMA over FP8, FP6 and FP4, as on MI450X, MI455X and MI430X.
capabilities.h:336 ↗
#define nk_cap_cdna5_k ((nk_capability_t)1 <<59)
nk_cap_metal_kApple baseline, which runs on every Metal device.
capabilities.h:339 ↗
#define nk_cap_metal_k ((nk_capability_t)1 <<60)
nk_cap_apple9_kApple GPU family 9, as on M3 and M4.
capabilities.h:342 ↗
#define nk_cap_apple9_k ((nk_capability_t)1 <<61)
nk_cap_apple10_kApple GPU family 10, as on M5.
capabilities.h:345 ↗
#define nk_cap_apple10_k ((nk_capability_t)1 <<62)
nk_cap_gpus_kEvery GPU capability above, which the CPU dispatch never detects, compiles or enables.
capabilities.h:348 ↗
#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_kEvery CPU capability, the bits below the first GPU vendor's.
capabilities.h:354 ↗
#define nk_cap_cpus_k (nk_cap_cuda_k - 1)
NUMKONG_CAPABILITIES_NAME_CAPACITYBuffer size nk_capabilities_name never overruns, including its null terminator.
capabilities.h:357 ↗
#define NUMKONG_CAPABILITIES_NAME_CAPACITY 1024