Initial Proof of Concept: Data Movement

Show forward progress on HiHAT

Demonstrate the layering of a simple app, the user layer, the common layer, glue code and implementation

One user-layer API, makes implementation-specific trade-offs to choose which common-layer API to call

Demonstrate low performance overheads

Retargetability for CPU in addition to GPU

What would be enabled for runtime clients

Basic data movement among multiple instances of multiple device kinds

Incremental goals

Illustrate interoperability with respect to user memory allocation

Introduce basic enumeration, e.g. for devices

Illustrate generality, e.g. with support for managed memory

Stretch goals

Demonstrate node to node transfers

Non-goals (at least initially)

Phase 1: simple data movement of HiHAT-allocated data

Choose API for pair of devices, based on address

Use allocator to associate address with device; introduces buffers

Assume that device kinds <CPU and GPU> and index order is <CPU0, GPU0, GPU1>

Assume that there's enough physical memory on host to pin all host-side memory

Batch with a mix of source and targets for GPU, CPU, no dependences

Same, but code using raw CUDA APIs instead of HiHAT

John Stone @ UIUC: molecular orbital algorithm from quantum chemistry (extracted from VMD)

alloc_buffer(void** out_address, size_t size, int device_kind_index, int device_index, int mem_kind)

move_data(void* target, void* src, size_t size), determine CPU vs. GPU based on address range, choose which common layer API based on source and target

move_{local_to_remote, remote_to_local, remote_to_remote}(void* dest_addr, void* src_addr, int dest_device_kind_index, int dest_device_index, int src_device_kind_index, int src_device_index, int src_device_id, size_t size)

move_local_to_local(void* dest_addr, void* src_addr, size_t size)

cudaMallocHost, cudaMalloc( void** address, sizte_t size )

cudaMemcpy(void* dest, const void* src, size_t size, cudaMemcpyKind cudaMemcpy{Host, Device}To{Host, Device})

memcpy(void* dest, void* src, size_t size)

Vary the payload size, show how % overhead from HiHAT APIs and glue code varies

Use cudaMemcpyAsync with a range of # of streams, as compared with blocking APIs only, for a range of batch sizes

Make a subset of them have a dependence chain, and compare all blocking vs. having dependent subsets in same streams

Compare cudaMemcpy(dest, src, size, cudaMemcpyDeviceToDevice) with cudaMemcpyPeer

Phase 2: simple data movement of non-HiHAT-allocated data

Enable wrapping existing memory in a buffer without allocating it in HiHAT

wrap_with_buffer(void* in_address, sizte_t size, int device_kind, int device_index)

Phase 3: basic device enumeration

Enumerate and specify devices

query_num_device_kinds(int *out_num_device_kinds, int PE_index)

query_device_kind(int *out_device_kind, int in_device_kind_index, int PE_index)

query_num_devices_of_kind(int *out_num_devices_of_kind, int in_device_kind_index, int PE_index)

Phase 4: basic managed memory

Demonstrate support for managed memory

alloc_buffer(void* out_address, size_t size, int device_kind_index=MANAGED, int device_index=MANAGED) (on this PE)

Choose cases that minimize overhead of managed memory

cudaMallocManaged

Compare managed vs. non-managed memory

Phase 5 (stretch): node to node transfers

Demonstrate node to node transfers

Presume that mpi_run is used to create multiple ranks

Default configuration uses MPI

init_communication() - inits MPI, sets up default communicator

query_num_PEs(int *out_num_PEs)

query_local_PE(int *out_PE_index)

Phase 6 (stretch): node to node transfers on selected transport

Select network transport

config_network_transport(<transport type>)

Vary the transport type, e.g. between MPI, UCX, TCP

Future directions

Xfers of multi-D and non-dense data, ability to use UCX

cudaMemAdvise for managed memory, cudaMemcpyToSymbol for constant memory

Ideas for lookups and trade-offs

Have to introduce notion of domains, need to use APIs to associate addresses with domains

Differentiating between frequent and infrequent and contriving motivation not to pin are a little contrived

CUDA handles this, doesn't offer a choice between DMA and aperture write