typedef void *(*) nk_allocate_t(nk_size_t bytes, void *handle, nk_stream_t stream);Every kernel takes plain pointers, so any CPU-accessible memory works: mmap, pinned buffers, CUDA unified memory, custom arenas. C++ views wrap any pointer without ownership. Owning containers accept any C++ Allocator.
template <typename T>
struct cuda_allocator {
using value_type = T;
T *allocate(std::size_t n) { T *p;
cudaMallocManaged(&p, n * sizeof(T), cudaMemAttachGlobal);
return p; }
void deallocate(T *p, std::size_t) noexcept { cudaFree(p); }
};
nk_dot_f32_best(cuda_managed_ptr, cuda_managed_ptr, 1024, &dot, capabilities, NULL); // C ABI, any pointer
auto view = nk::tensor_view<nk::f32_t>(mmap_ptr, rows, columns); // non-owning view
auto v = nk::vector<float, cuda_allocator<float>>::zeros(1024); // allocator-aware owningAPI
Types
nk_allocate_tAllocation callbacks retain the caller's context and use its device stream.
nk_allocator_tAllocation policy; the caller owns the handle and any stream passed to its callbacks.
Functions
nk_allocator_init_heapHeap allocations with power-of-two alignment; invalid alignment preserves the policy.
nk_status_t nk_allocator_init_heap(
nk_allocator_t *allocator,
nk_size_t alignment);nk_allocator_init_arenaA fixed-capacity arena; invalid storage or alignment preserves both the policy and buffer.
nk_status_t nk_allocator_init_arena(
nk_allocator_t *allocator,
void *buffer,
nk_size_t bytes,
nk_size_t alignment);nk_allocator_init_unified_bestSelects host/device shared storage; failures preserve the allocator.
nk_status_t nk_allocator_init_unified_best(
nk_allocator_t *allocator,
nk_capability_t capabilities);Backends
nk_allocator_init_unifiedserial
nk_allocator_init_device_bestSelects device-local storage, or shared storage on Metal.
nk_status_t nk_allocator_init_device_best(
nk_allocator_t *allocator,
nk_capability_t capabilities);CPU masks have no device allocator.
nk_allocator_init_pinned_bestSelects pinned host storage on CUDA/ROCm or shared storage on Metal; unsupported on CPU.
nk_status_t nk_allocator_init_pinned_best(
nk_allocator_t *allocator,
nk_capability_t capabilities);nk_memory_allocate_unified_bestAllocates bytes that the host and the device of stream both address.
nk_status_t nk_memory_allocate_unified_best(
nk_size_t bytes,
void **pointer,
nk_capability_t capabilities,
nk_stream_t stream);Backends
nk_memory_allocate_unifiedserial
CUDA and ROCm hand out managed memory, CUDA from the device's managed pool in the order of stream, and Metal a shared buffer every kernel on that device binds, so the host reads what a kernel wrote once the stream is synchronized.
Parameters
bytes | in | The size of the block, where zero hands out a null one. |
|---|---|---|
pointer | out | The block, or null on failure. CPU storage uses nk_default_alignment_k alignment. |
capabilities | in | The group to allocate for, like nk_cuda_capabilities_enabled reports. |
stream | in | Null on the CPU, or a stream of that group, naming the device. |
Statuses
nk_missing_gpu_kA group without a device or missing from this build.
nk_memory_free_unified_bestReturns a block of nk_memory_allocate_unified_best once the work queued on stream is done with it.
nk_status_t nk_memory_free_unified_best(
void *pointer,
nk_size_t bytes,
nk_capability_t capabilities,
nk_stream_t stream);Backends
nk_memory_free_unifiedserial
CUDA frees blocks of its managed pool once the work queued on stream completes, without waiting; devices without concurrent managed access and ROCm wait for the device first; Metal releases the buffer once the work committed to stream so far completes.
Parameters
pointer | in | The block, or null for none. |
|---|---|---|
bytes | in | The size it was allocated with. |
capabilities | in | One device's capabilities, like nk_cpu_capabilities_enabled reports. |
stream | in | The stream the block was allocated on, or null for the default. |
Statuses
nk_device_memory_mismatch_kA block the group never handed out.nk_missing_gpu_k
Macros
NUMKONG_DEFAULT_ALIGNMENTDefault allocation alignment, conservatively sized for cache interference.
#define NUMKONG_DEFAULT_ALIGNMENT 128