Quick Start: Data

Container

Category

There are three fundamental containers in ZPC:

• Vector<T, Allocator>: an array of elements of type T stored in a contiguous chunk of memory. T should be default constructible and trivially copyable.
• HashTable<Tn, dim, Index>: a hash table that maps coordinates (type: vec<Tn, dim>) to a one-dimension index (type: Index, e.g. int) starting from 0.
• bht<Tn, dim, Index, B, Allocator>: A better implemented bucketed hash table (in terms of load factor and hash performance) which can be used as a substitute for the above one. B is the bucket size.
• TileVector<T, tile_length, Allocator>: an aosoa-layout tile array, where each tile contains tile_length elements that have a certain amount of properties represented by N channels (specified at runtime) of type T (default constructible and trivially copyable).

Their underlying storage could be from one of the following memory spaces: host (memsrc_e::host), device (memsrc_e::device), unified memory (memsrc_e::um), etc. All memory spaces should implement the following set of memory operations, including:

void *allocate(tag<memsrc_e>, size_t size, size_t alignment)
void deallocate(tag<memsrc_e>, void *ptr, size_t size, size_t alignment)
void copy(tag<memsrc_e>, void *dst, void *src, size_t size). ps. the interface might change in the future
void memset(tag<memsrc_e>, void *ptr, int chval, size_t size)
void advise(tag<memsrc_e>, std::string_view advice, void *addr, size_t size, ProcID did)

In ZPC, a memory location consists of a memory space and a device index:

struct MemoryLocation {
  memsrc_e _memsrc;
  ProcID _devid; // processor id
};
// following structs are implemented upon MemoryLocation
struct MemoryProperty {
  MemoryLocation location;
  MemoryTraits traits;
}; // aliased as MemoryHandle for legacy usage..
struct MemoryEntity {
  MemoryLocation location;
  void *ptr;
};

Note that the device index for the host-side main memory is -1, whereas the index for GPU devices starts from 0.

The aforementioned three types of containers are all move-(assignment)-constructible and copy-(assignment)-constructible, and additionally provide a special function clone(MemoryLocation handle), which clones (deep-copy) all its content to the memory location specified by location (might be different from the original one).

Construct

Vector<T> a{memsrc_e::um, -1};  ///< "default" ctor, stored on host
Vector<T> b{10, memsrc_e::um, 0};  ///< allocate an array of 10 elements within unified memory, preferred storage location is 0-th GPU 
Vector<T> c{b};  ///< a copy from b
Vector<T> d{std::move(b)};  ///< moved from b
Vector<T> e = b.clone({memsrc_e::host, -1});

The constructors for HashTable is similar to the Vector container. TileVector, however, needs programmers to specify a property list in the first argument, e.g. TileVector<T, 32> tv{{{"density", 1}, {"pos", 3}}, N, memsrc_e::host, -1}.

Acceleration/Spatial Data Structures

ZPC also supports a range of acceleration/spatial data structures that are needed for efficient physics-based simulations.

• LBvh<dim, tile_length, T, Allocator>: Linear BVH for fast spatial object query acceleration.
• SparseGrid<dim, T, SideLength, Allocator, Ti>: Sparse uniform grid data structure.
• VdbGrid<dim, T, BitSeq, Allocator> (WIP): Adaptive grid data structure (openvdb-alike).
• AdaptiveTileTree<dim, T, NumLevels, N, Allocator> (WIP): Adaptive tile tree data structure (bifrost-alike). B+ tree as an alternative acceleration structure is also intriguing, we may implement it in the future.

Headers

#include "zensim/container/..."
#include "zensim/geometry/..."

Access in a Compute Backend

The above containers and spatial data structures are not directly accessible in a hand-written function for a certain compute back-end (e.g. openmp, cuda for now) specified by a execspace_e tag. They should all be wrapped within a proxy object (can also be interpreted as a view) before being used in a kernel.

__global__ void display_vector(VectorProxy<execspace_e::cuda, Vector<float>> vec) {
  printf("%d: %f\n", (int)threadIdx.x, vec(threadIdx.x));
}

int main() {
  Vector<float> a{10, memsrc_e::um, 1};
  // init a
  // display a
  display_vector<<<1, 10>>>(proxy<execspace_e::cuda>(a));
  return 0;
}