float * of C/C++/Rust Fails to Distinguish Host DRAM from GPU HBM

Aditya Kumar · 23 September 2026

TL;DR — malloc and cudaMalloc return memory that lives in different places. They are supposed to be read by different processors, and run out at different sizes. C, C++ and Rust give both results the same type. Vx makes the memory space (placement) part of the type, so a value from device memory cannot be passed, assigned or read as if it were host memory.

An allocator knows exactly one thing better than anyone else in the program: which memory it just handed out. The first thing the language does with that answer is throw it away.

Here is a simple example in C to demonstrate the problem

float *a = (float *)malloc(16 * sizeof(float));
float *b;
cudaMalloc((void **)&b, 16 * sizeof(float));

a = b;              // compiles
float x = b[0];    // compiles; faults on a discrete GPU

a points into host DRAM. b points into the GPU's HBM, which a host thread on a discrete GPU cannot address at all. They are the same type, they assign to each other, and reading through either is a well-formed expression.

Look at where the difference is written down. The signature is

cudaError_t cudaMalloc(void **devPtr, size_t size);

The only record that the result is device memory is the function's name and the parameter name devPtr. Both are comments, as far as the compiler is concerned.

C++ and Rust agree with C

C++ inherits the same float *. CUDA's __device__ and __shared__ qualify a variable's declaration; they do not survive into a pointer's type, and inside a kernel a pointer is generic until something proves otherwise. Rust has one raw pointer per mutability, and its foreign-function declarations come out identical:

extern "C" {
    fn malloc(size: usize) -> *mut c_void;
    fn cudaMalloc(dev_ptr: *mut *mut c_void, size: usize) -> i32;
}

Rust's memory model does track more about a pointer than its address. It tracks provenance: which allocation the pointer was derived from. It does not track which memory that allocation is in, because in its abstract machine there is only one.

Popular libraries have made efforts to address this

People who write GPU code hit this early, so the libraries built the missing type themselves. Thrust has thrust::device_ptr<T>. SYCL has multi_ptr<T, address_space>. In Rust, cudarc has CudaSlice<T> and cust has DeviceBuffer<T>. Inside one of these libraries, a device buffer and a host buffer are different types.

Each one is their own vocabulary, though, and they meet at the raw pointer. cuBLAS takes const float *. A kernel launch takes float *. Thrust's way out is thrust::raw_pointer_cast, and the Rust crates have their own equivalents. Every time a buffer crosses from one library to another, it crosses as a bare pointer, and the space is gone again.

The compiler underneath them knows the fact matters. LLVM's IR has address spaces (ptr addrspace(1) is NVPTX global memory, addrspace(3) is shared), and its NVPTX pipeline runs a pass called InferAddressSpaces whose job is to work out, after the fact, which space each generic pointer points into. The allocator knew. The type dropped it. A compiler pass spends its time guessing it back, and it gives up at any call it cannot see through.

The bugs highlight the underlying problem

This is a mistake people make often enough to leave a trail. A few from public trackers:

The cuBLAS row is the clearest. The library's authors knew exactly which memory a result pointer names, and the only place they could record it was a mode flag on a handle, because the parameter's type has nowhere to put it. The last row is a compiler making the same mistake a person does.

Prior art in programming languages

Putting the space in the pointer's type is an old idea. OpenCL C writes __global float * and __local float * as distinct types. ISO/IEC TR 18037 gave C named address spaces for embedded targets, and clang accepts __attribute__((address_space(N))) in C and C++. Sequoia, Chapel's locales, X10's places and Legion's regions all built placement into the language.

CUDA went the other way on purpose. Unified virtual addressing, since CUDA 4.0, gives every allocation a distinct address across host and device, so one generic pointer can reach any memory and cudaPointerGetAttributes reports at run time where it points. That buys one kernel for every memory, and it moves the question of where a pointer points from the compiler to the run time.

Vx keeps the space in the type and adds a description of the machine to check it against: which spaces each processor can see, how large each space is, and which links join them. The check runs at compile time, and the error names the transfer that fixes it.

In Vx, the space is in the type

A Vx tensor's type has a slot for the memory it lives in. Here is a host function, and a caller that passes it device memory:

fn sum_on_host(t : Tensor<f32, [4, 4]>) -> f32 {
  return t[0][0] + t[1][1];
}

fn main() -> i32 {
  let host_data : Tensor<f32, [4, 4]> = Tensor<f32, [4, 4]>::new();
  let on_gpu = transfer(host_data, Memory::GPU_HBM);
  print(sum_on_host(on_gpu));
  return 0;
}

The compiler refuses the call twice, once for the type and once for the read:

Error[E6003] at 8:21: 'on_gpu' lives in GPU_HBM but CPU sees only [CPU_DRAM, NPU_HBM];
  insert an explicit transfer to CPU_DRAM (cost 50 on the declared path)
Error[E3003] at 8:9: Type mismatch in argument 1 for function 'sum_on_host'.
  Expected Tensor<f32, [4, 4]>, got Tensor<f32, [4, 4], GPU, GpuHbm>

The other ways C lets the two mix are closed too. Assigning device memory to a variable declared as a host tensor is E3001. Reading it inside an unsafe block is still E6003, because unsafe in Vx permits unchecked operations, and reading memory your processor cannot address is simply wrong. Taking the element's address to get a raw pointer is refused at the same read.

The fix is the copy the C program also needed:

fn sum_on_host(t : Tensor<f32, [4, 4]>) -> f32 {
  return t[0][0] + t[1][1];
}

fn main() -> i32 {
  let host_data : Tensor<f32, [4, 4]> = Tensor<f32, [4, 4]>::new();
  let on_gpu = transfer(host_data, Memory::GPU_HBM);
  let home = transfer(on_gpu, Memory::CPU_DRAM);
  print(sum_on_host(home));
  return 0;
}

That compiles and runs. Nothing about it needs a GPU on the machine doing the compiling: which spaces a CPU can see comes from the machine description the compiler checks against.

Where the space goes in a unified memory system

On Apple silicon the CPU and the neural engine share physical memory, and the built-in CPU description lists NPU_HBM among the spaces it can see. A host read of a tensor in NPU_HBM is legal there, and Vx accepts it. The transfer into that space is still written, and compiles to almost nothing. The type records where a value lives; the machine description decides who may read it.

Hardware is moving this way. On Grace Hopper the CPU reads the GPU's HBM directly over NVLink-C2C, and MI300A puts CPU and GPU on one pool of HBM. Where a read is legal it still has a price: a remote read runs at the link's bandwidth, and on a two-socket host a remote NUMA access measured 2.35× the cost of a local one. The machine file records both facts: visible: says who may read a space, and the transfer edges say what moving data costs.

Even in Vx raw pointers mean CPU memory

Vx has raw pointers too, *const T and *mut T, and they carry no space parameter. That is deliberate: a raw pointer points into the default CPU memory, which is what a float * means to C, C++ and Rust code today. A pointer passed to or from a C library keeps the meaning it had there, so existing code and extern declarations work unchanged.

Memory anywhere else is reached through a type that names its space: Tensor or Pinned. That is why taking the address of a device element from host code is refused. The pointer it would produce claims CPU memory, and the element is in HBM.

That leaves a question for GPU libraries: how does a device tensor reach cuBLAS? Today the runtime does it. Its CUDA dispatch calls cuBLAS for matrix products, with the buffers it already holds on the device. User code cannot yet pass a device tensor to a C library through its own extern declaration. That needs a way to spell a device-memory parameter in an extern signature, and Vx does not have one yet; it is tracked in issue #742.

An earlier post, What a pointer forgets, follows the same fact across a function return. The heterogeneous model covers placement and transfer, and E6003 has the diagnostic's own page.