Halloc GPU memory allocator, version 0.11

INTRO

Halloc is a high-throughput malloc/free-style dynamic memory allocator for
NVidia Kepler GPUs. It is based on using bit arrays to represent free blocks and
using a hash function to quickly search for free blocks. This idea, combined
with clever slab management and performance tuning, enables a really fast
allocator. Halloc achieves more than 1.5 bln. mallocs/s (more than 1
bln. malloc/free pairs/s) on K20X and 16-byte allocations, with tens of
thousands of GPU threads and more than 100MiB allocated. This is much higher
than other state-of-the-art GPU allocators. In addition, halloc's performance is
also more stable. This makes halloc suitable for use in GPGPU applications
requiring fast dynamic memory management. Halloc is mainly designed for small
allocation sizes, and delegates allocations larger than 3KiB to CUDA allocator.


REQUIREMENTS

Software: CUDA 5.0 or higher (tested with 6.5)
Hardware: Compute Capability 2.0 or higher (tested on CC 3.5 devices K20X and K40).

Note: libraries and tests are currently not compiled for compute_50/sm_50, i.e. Maxwell.


COMPILING

To compile halloc library, type (in project's top directory):

	make

To run correctness tests (CAUTION: takes a lot of time!):

	make test

To build correctness tests without running them:

	make build-corr

To build performance tests without running them:
	 
	make build-perf

Performance tests are then located in ./tst/perf/bin directory, and can be
invoked individually, e.g.

	./tst/perf/bin/throughput
	./tst/perf/bin/phase-throughput -f0.95 -F0.05 -e0.91 -g5 -t128

To install, edit PREFIX variable in the makefile to your desired install
directory (default ~/usr) and type:

	make install

To uninstall:

	make uninstall


USING HALLOC 

See samples/ directory for samples using Halloc.

Compiling Your Program

The GPU application then needs to be compiled with halloc static library using
separate device compilation and linking. Assuming that the variable $PREFIX
contains the installation prefix, and myprog.cu is the file being compiled, this
can be done as follows:

  nvcc -arch=sm_35 -O3 -I $(PREFIX)/include -dc myprog.cu -o myprog.o
	nvcc -arch=sm_35 -O3 -L $(PREFIX)/lib -lhalloc -o myprog myprog.o


Halloc API 

The functions defined by halloc are in the halloc.h file, which needs to be
included into your code to use halloc:

#include <halloc.h>

Before using halloc, in device functions it has to be initialized with ha_init()
function:

void ha_init(halloc_opts_t opts = halloc_opts_t());

It can be given a full halloc_opts_t structure to control fine halloc
parameters, such as slab size or fraction of used chunks at which the slab is
considered "busy". It can also be called just with specifying amount of memory
to allocate, or completely without any parameter list to preserve defaults:

ha_init(512 * 1024 * 1024);  // pass memory to allocate
ha_init();  // use default amount of memory

Halloc defines two functions, hamalloc to allocate and hafree to free memory
(malloc and free are used by CUDA allocator, therefore halloc has to use other
names). These functions can only be called from device code.

void *hamalloc(size_t nbytes);
void hafree(void *p);

Otherwise, these functions have pretty much the same behavior as standard C
malloc/free, e.g.:

// allocate an array
int *p = (int *)hamalloc(8 * sizeof(int));
p[0] = 0;
p[1] = threadIdx.x;
p[2] = 2;
// ...
// free the array
hafree(p);


// allocate a list
typedef struct list_ {
	int element;
	struct list_ *next;
} list;
// ...
list *l = (list *)hamalloc(sizeof(list));
l->element = 1;
l->next = (list *)hamalloc(sizeof(list));
l->next->element = 2;
l->next->next = NULL;

The functions can be used in pretty much the same way as in C programs. hamalloc
accepts the number of bytes to allocate, and returns the pointer to allocated
memory, or NULL if memory cannot be allocated. Similarly, hafree accepts either
a pointer returned by hamalloc or NULL, and frees the memory previously
allocated. Naturally, hamalloc and hafree are thread-safe, and can be called
simultaneously by threads of the same or different kernels. hamalloc allocations
persist across kernel invocations, and can be used in other kernel
calls. Pointers allocated by hamalloc can only be freed by hafree;
they cannot be deallocated, e.g., by host/device cudaFree/free.

ha_shutdown() is intended to free resources used by halloc, but is currently a
no-op.


LIMITATIONS

There is currently no way to change parameters or allocate more memory after
halloc has been initialized.


BUGS

Though correctness tests pass successfully, this provies nothing, of
course. Some bugs are most likely there ;)