# External Memory

> NumKong for C

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.

```cpp
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 owning
```

`#include <numkong/memory.h>`

`#include <numkong/types.h>`

## nk_allocate_t

Allocation callbacks retain the caller's context and use its device stream.

```c
typedef void *(*) nk_allocate_t(nk_size_t bytes, void *handle, nk_stream_t stream);
```

## nk_allocator_t

Allocation policy; the caller owns the handle and any stream passed to its callbacks.

```c
struct nk_allocator_t
```

- `allocate` `nk_allocate_t`
- `free` `nk_free_t`
- `handle` `void *`

## nk_allocator_init_heap

Heap allocations with power-of-two `alignment`; invalid alignment preserves the policy.

```c
nk_status_t nk_allocator_init_heap(
    nk_allocator_t *allocator,
    nk_size_t alignment);
```

## nk_allocator_init_arena

A fixed-capacity arena; invalid storage or alignment preserves both the policy and buffer.

```c
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_best

Selects host/device shared storage; failures preserve the allocator.

```c
nk_status_t nk_allocator_init_unified_best(
    nk_allocator_t *allocator,
    nk_capability_t capabilities);
```

- `nk_allocator_init_unified` backends: serial

## nk_allocator_init_device_best

Selects device-local storage, or shared storage on Metal.

```c
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_best

Selects pinned host storage on CUDA/ROCm or shared storage on Metal; unsupported on CPU.

```c
nk_status_t nk_allocator_init_pinned_best(
    nk_allocator_t *allocator,
    nk_capability_t capabilities);
```

## nk_memory_allocate_unified_best

Allocates `bytes` that the host and the device of `stream` both address.

```c
nk_status_t nk_memory_allocate_unified_best(
    nk_size_t bytes,
    void **pointer,
    nk_capability_t capabilities,
    nk_stream_t stream);
```

- `nk_memory_allocate_unified` backends: serial

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.

- `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`](https://www.unum.cloud/docs/numkong/c/runtime-dispatch.md#nk_cuda_capabilities_enabled) reports.
- `stream` (in): Null on the CPU, or a stream of that group, naming the device.

Statuses: - [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k) A group without a device or missing from this build.

## nk_memory_free_unified_best

Returns a block of [nk_memory_allocate_unified_best](#nk_memory_allocate_unified_best) once the work queued on `stream` is done with it.

```c
nk_status_t nk_memory_free_unified_best(
    void *pointer,
    nk_size_t bytes,
    nk_capability_t capabilities,
    nk_stream_t stream);
```

- `nk_memory_free_unified` backends: serial

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.

- `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`](https://www.unum.cloud/docs/numkong/c/runtime-dispatch.md#nk_cpu_capabilities_enabled) reports.
- `stream` (in): The stream the block was allocated on, or null for the default.

Statuses: - [`nk_device_memory_mismatch_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_device_memory_mismatch_k) A block the group never handed out.
- [`nk_missing_gpu_k`](https://www.unum.cloud/docs/numkong/c/quickstart.md#nk_missing_gpu_k)

## NUMKONG_DEFAULT_ALIGNMENT

Default allocation alignment, conservatively sized for cache interference.

```c
#define NUMKONG_DEFAULT_ALIGNMENT 128
```