diff --git a/Makefile b/Makefile
new file mode 100644
index 0000000..da0c8f2
--- /dev/null
+++ b/Makefile
@@ -0,0 +1,19 @@
+NAME = test_ompcl
+
+SRCS = $(shell find . -maxdepth 1 -type f -name '*.c')
+OBJS = $(SRCS:.c=.o)
+
+LDLIBS += -lOpenCL
+
+.PHONY: all clean test
+
+all: ${NAME}
+
+${NAME}: ${OBJS}
+	${CC} ${LDFLAGS} $^ ${LDLIBS} -o $@
+
+test:
+	./${NAME}
+
+clean: 
+	@${RM} ${NAME} ${OBJS}
diff --git a/ccubes.cl b/ccubes.cl
new file mode 100644
index 0000000..05e1a14
--- /dev/null
+++ b/ccubes.cl
@@ -0,0 +1,286 @@
+#ifdef USE_DOUBLE
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+typedef double real;
+#define R_ZERO	1e-14
+#else
+typedef float real;
+#define R_ZERO	1e-10
+#endif
+
+#define BITS_PER_WORD 32
+
+#define ROW_DIM 0
+#define COL_DIM 1
+
+// #pragma OPENCL EXTENSION cl_amd_printf : enable
+// #pragma OPENCL EXTENSION cl_khr_select_fprounding_mode : enable
+// #pragma OPENCL SELECT_ROUNDING_MODE rtz
+
+#ifdef RANGE_DEBUG
+#define RANGE_CHECK(lower, upper, value, str) do {	\
+	if ((value) < (lower) || (value) > (upper)) {	\
+		printf("%s", (str));			\
+		return;					\
+	}						\
+} while(0);
+#else
+#define RANGE_CHECK(lower, upper, value, str)
+#endif
+
+unsigned long int
+nchoosek(int n, int k)
+{
+    if (k == 0 || k == n) return 1;
+    if (k == 1) return n;
+
+    unsigned long int result = 1;
+
+    if (k > n - k) {
+        k = n - k;
+    }
+
+    for (int i = 0; i < k; i++) {
+        result = result * (n - i) / (i + 1);
+    }
+
+    return result;
+}
+
+/*
+ *
+ * PROBLEM: CCubes
+ *
+ * INPUT:
+ *	k - current input
+ *	ninputs - number of inputs
+ *	posrows - positive output rows (the ON set)
+ *	negrows - negative output rows (the OFF set)
+ *	pichart_words - words needed per PI chart columns
+ *	implicant_words - words needed per PI representation
+ *	nofvalues (ninputs x 1) - number of values
+ *	nofpi (ninputs x 1) - number of prime implicants
+ *	ON_set (posrows x ninputs) - ON set
+ *	OFF_set (ninputs x negrows) - OFF set
+ *
+ * OUTPUT:
+ *	x (n x 1) - solution (L \ b)
+ *
+ * NOTE: Both input and output must be allocated before calling this funciton.
+ */
+__kernel void
+ccubes_task(int k, int ninputs,
+		int posrows,
+		int negrows,
+		int pichart_words,
+		int implicant_words,
+		__global const real *nofvalues,
+		__global const real *nofpi,
+		__global const real *ON_set,
+		__global const real *OFF_set,
+		__global const unsigned int *p_implicants_pos,
+		__global const unsigned int *p_implicants_val,
+		__global const int *last_index,
+		__global const int *p_covered,
+		__global const int *p_pichart_pos,
+		)
+{
+	/* work-item?: task in nchoosek(ninputs, k) */
+	/* work-group?: k in 1 to ninputs */
+	/* total work: tasks in nchoosek for k in 1 to ninputs */
+
+	size_t task = get_global_id(0);
+	
+	int prevfoundPI = 0;
+
+	int tempk[k]; /* max is tempk[ninputs] */
+	int x = 0;
+	int start_point = task;
+
+	// fill the combination for the current task
+	for (int i = 0; i < k; i++) {
+		while (nchoosek(ninputs - (x + 1), k - (i + 1)) <= start_point) {
+			start_point -= nchoosek(ninputs - (x + 1), k - (i + 1));
+			x++;
+		}
+		tempk[i] = x;
+		x++;
+	}
+
+	// allocate vectors of decimal row numbers for the positive and negative rows
+	int decpos[posrows];
+	int decneg[negrows];
+
+	// create the vector of multiple bases, useful when calculating the decimal representation
+	// of a particular combination of columns, for each row
+	int mbase[k];
+	mbase[0] = 1; // the first number is _always_ equal to 1, irrespective of the number of values in a certain input
+
+	// calculate the vector of multiple bases, for example if we have k = 3 (three inputs) with
+	// 2, 3 and 2 values then mbase will be [1, 2, 6] from: 1, 1 * 2 = 2, 2 * 3 = 6
+	for (int i = 1; i < k; i++) {
+		mbase[i] = mbase[i - 1] * nofvalues[tempk[i - 1]];
+	}
+
+	// calculate decimal numbers, using mbase, fills in decpos and decneg
+	for (int r = 0; r < posrows; r++) {
+		decpos[r] = 0;
+		for (int c = 0; c < k; c++) {
+			decpos[r] += ON_set[tempk[c] * posrows + r] * mbase[c];
+		}
+	}
+
+	for (int r = 0; r < negrows; r++) {
+		decneg[r] = 0;
+		for (int c = 0; c < k; c++) {
+			decneg[r] += OFF_set[tempk[c] * negrows + r] * mbase[c];
+		}
+	}
+
+
+	int possible_rows[posrows];
+
+	bool possible_cover[posrows];
+	possible_cover[0] = true; // bool flag, to be set with false if found among the OFF set
+
+	int found = 0;
+
+	// identifies all unique decimal rows, for the selected combination of k inputs
+	for (int r = 0; r < posrows; r++) {
+		int prev = 0;
+		bool unique = true; // bool flag, assume the row is unique
+		while (prev < found && unique) {
+			unique = decpos[possible_rows[prev]] != decpos[r];
+			prev++;
+		}
+
+		if (unique) {
+			possible_rows[found] = r;
+			possible_cover[found] = true;
+			found++;
+		}
+	}
+
+	if (found > 0) {
+		// some of the ON set numbers are possible PIs (not found in the OFF set)
+		int frows[found];
+
+		// verify if this is a possible PI
+		// (if the same decimal number is not found in the OFF set)
+		for (int i = found - 1; i >= 0; i--) {
+			int j = 0;
+			while (j < negrows && possible_cover[i]) {
+				if (decpos[possible_rows[i]] == decneg[j]) {
+					possible_cover[i] = false;
+					found--;
+				}
+				j++;
+			}
+
+			if (possible_cover[i]) {
+				frows[found - i - 1] = possible_rows[i];
+			}
+		}
+		// Rprintf("task: %d; rows: %d\n", task, found);
+
+		for (int f = 0; f < found; f++) {
+
+
+			// create a temporary vector of length k, containing the values from the initial ON set
+			// plus 1 (because 0 now signals a minimization, it becomes 1, and 1 becomes 2 etc.
+			int tempc[k];
+
+			// using bit shifting, store the fixed bits and value bits
+			unsigned int fixed_bits[implicant_words];
+			unsigned int value_bits[implicant_words];
+
+			for (int i = 0; i < implicant_words; i++) {
+				fixed_bits[i] = 0U;
+				value_bits[i] = 0U;
+			}
+
+			for (int c = 0; c < k; c++) {
+				int value = ON_set[tempk[c] * posrows + frows[f]];
+				tempc[c] = value + 1;
+
+				int word_index = tempk[c] / BITS_PER_WORD;
+				int bit_index = tempk[c] % BITS_PER_WORD;
+
+				fixed_bits[word_index] |= 1U << bit_index;
+				value_bits[word_index] |= (unsigned int)value << (bit_index * value_bit_width);
+			}
+
+			// check if the current PI is not redundant
+			bool redundant = false;
+
+			int i = 0;
+			while (i < prevfoundPI && !redundant) {
+				// /*
+				// - ck contains the complexity level for each of the previously found non-redundant PIs
+				// - indx is a matrix containing the indexes of the columns where the values were stored
+				// - a redundant PI is one for which all values from a previous PI are exactly the same:
+				// 0 0 1 2 0, let's say previously found PI
+				// which means a corresponding ck = 2 and a corresponding indx = [3, 4]
+				// 0 0 1 2 1 is redundant because on both columns 3 and 4 the values are equal
+				// therefore sumeq = 2 and it will be equal to v = 2 when reaching the complexity level ck = 2
+				// */
+
+				bool is_subset = true; // Assume it's a subset unless proven otherwise
+
+				for (int w = 0; w < implicant_words; w++) {
+					// If the new PI has values on positions outside the existing PI’s fixed positions, it’s not a subset
+					if ((fixed_bits[w] & p_implicants_pos[i * implicant_words + w]) != p_implicants_pos[i * implicant_words + w]) {
+						is_subset = false;
+						break;
+					}
+
+					// then compare the value bits, if one or more values on those positions are different, it’s not a subset
+					if ((value_bits[w] & p_implicants_val[i * implicant_words + w]) != p_implicants_val[i * implicant_words + w]) {
+						is_subset = false;
+						break;
+					}
+				}
+
+				redundant = is_subset;
+
+				i++;
+			}
+
+			if (redundant) continue;
+
+			bool coverage[posrows];
+			int covsum = 0;
+			unsigned int pichart_values[pichart_words];
+			for (int w = 0; w < pichart_words; w++) {
+				pichart_values[w] = 0U;
+			}
+
+			for (int r = 0; r < posrows; r++) {
+				coverage[r] = decpos[r] == decpos[frows[f]];
+				if (coverage[r]) {
+					int word_index = r / BITS_PER_WORD;
+					int bit_index = r % BITS_PER_WORD;
+					pichart_values[word_index] |= (1U << bit_index);
+				}
+				covsum += coverage[r];
+			}
+
+			// verify row dominance
+			int rd = 0;
+			while (rd < last_index[covsum - 1] && !redundant) {
+
+				bool dominated = true;
+				for (int w = 0; w < pichart_words; w++) {
+					if ((pichart_values[w] & p_pichart_pos[p_covered[rd] * pichart_words + w]) != pichart_values[w]) {
+						dominated = false;
+						break;
+					}
+				}
+
+				redundant = dominated;
+				rd++;
+			}
+
+			if (redundant) continue;
+		}
+	}
+}
diff --git a/cl_setup.c b/cl_setup.c
new file mode 100644
index 0000000..58ccee7
--- /dev/null
+++ b/cl_setup.c
@@ -0,0 +1,373 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include "logging.h"
+
+#include "cl_setup.h"
+#define CL_DEBUG
+
+cl_int
+cl_init(struct cl_uctx *puctx)
+{
+	/* OpenCL specific variables */
+	size_t dataBytes;
+
+	cl_int result = CL_SUCCESS;
+	cl_uint nplat = 0, ndevices = 0, i = 0, matched_plat;
+	cl_platform_id *platforms = NULL, platform = NULL;
+	cl_device_id *devices = NULL;
+	cl_context_properties props[3] = {CL_CONTEXT_PLATFORM, 0, 0};
+	cl_context ctx = NULL;
+
+
+	if (puctx == NULL || puctx->platform_name == NULL) {
+		return CL_INVALID_VALUE;
+	}
+	if (puctx->device_type > CL_DEVICE_TYPE_GPU) {
+		puctx->device_type = CL_DEVICE_TYPE_ALL;
+	}
+
+	/*
+	 * Initialize OpenCL.
+	 */
+	result = clGetPlatformIDs(0, NULL, &nplat);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "Failed getting the number of platforms");
+		goto err;
+	}
+	if (nplat < 0) {
+		log_warn("cl", "No platforms found");
+		goto err;
+	}
+	platforms = calloc(nplat, sizeof platforms[0]);
+	if (platforms == NULL) {
+		log_warn("cl", "Failed to allocate platforms");
+		goto err;
+	}
+	result = clGetPlatformIDs(nplat, platforms, NULL);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "Failed fetching the platforms");
+		goto err;
+	}
+
+	log_debug("cl", "Found %d platforms", nplat);
+	for (i = 0; i < nplat; i++) {
+		char platname[100];
+		result = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR,
+		    sizeof(platname), platname, NULL);
+		if (result != CL_SUCCESS) {
+			log_warn("cl", "Failed fetching platform info");
+			goto err;
+		}
+		log_debug("cl", "Checking %s == %s",
+		    puctx->platform_name, platname);
+		if (!strcmp(platname, puctx->platform_name)) {
+			platform = platforms[i];
+			matched_plat = i;
+			break;
+		}
+	}
+	if (platform == NULL) {
+		log_warn("cl", "No matching platform found");
+		result = CL_DEVICE_NOT_FOUND;
+		goto err;
+	}
+
+#ifdef CL_DEBUG
+	printf("-----------------------------------------------------------\n");
+	printf("   PLATFORM   INFORMATION (the number of platforms = %d)   \n",
+	    nplat);
+	for(i = 0 ; i < nplat; i++)
+	{
+		char *long_str;
+		char str[1024];
+		size_t str_size;
+
+		printf("-------------------------------------------------------"
+		    "----\n");
+		printf( "    PLATFORM ID : %d                                  "
+		    "\n", i);
+		printf("-------------------------------------------------------"
+		    "----\n");
+
+		clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, sizeof(str),
+		    str, NULL);
+		printf("Platform name    : %s\n", str);
+
+		clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION,
+		    sizeof(str), str, NULL);
+		printf("Platform version : %s\n", str);
+
+		clGetPlatformInfo(platforms[i], CL_PLATFORM_EXTENSIONS, 0,
+		    NULL, &str_size);
+		long_str = (char *)malloc(str_size);
+		clGetPlatformInfo(platforms[i], CL_PLATFORM_EXTENSIONS,
+		    str_size, long_str, NULL);
+		printf("Platform extensions : %s\n", long_str);
+
+		printf("-------------------------------------------------------"
+		    "----\n\n");
+
+		free(long_str);
+	}
+#endif
+
+	props[1] = (cl_context_properties)platform;
+
+	ctx = clCreateContextFromType(props, CL_DEVICE_TYPE_ALL, NULL, NULL,
+	    &result);
+	if(result != CL_SUCCESS) {
+		log_warn("cl", "Failed to create context");
+		goto err;
+	}
+
+	/* fetch the list of devices associated with context */
+	result = clGetContextInfo(ctx, CL_CONTEXT_DEVICES, 0, NULL,
+		&dataBytes);
+	if(result != CL_SUCCESS) {
+		log_warn("cl", "Failed to fetch devices size!");
+		goto err;
+	}
+	devices = (cl_device_id *)malloc(dataBytes);
+	if (devices == NULL) {
+		log_warn("cl", "devices malloc() failed!");
+		goto err;
+	}
+	result |= clGetContextInfo(ctx, CL_CONTEXT_DEVICES, dataBytes,
+	    devices, NULL);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "clGetContextInfo() failed with %d!", result);
+		goto err;
+	}
+
+
+#ifdef CL_DEBUG
+	result = clGetDeviceIDs(platforms[matched_plat], CL_DEVICE_TYPE_ALL,
+	    0, NULL, &ndevices);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "clGetDeviceIDs() failed!");
+		goto err;
+	}
+	printf("-----------------------------------------------------------\n");
+	printf("      DEVICE   INFORMATION (the number of devices = %d)    \n",
+	    ndevices);
+	for(i = 0 ; i < ndevices ; i++)
+	{
+		char    str[1024];
+		size_t  int_info;
+
+		printf("-------------------------------------------------------"
+		    "----\n");
+		printf("    DEVICE ID : %d                                     "
+		    "\n",i);
+		printf("-------------------------------------------------------"
+		    "----\n");
+
+		clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(str), str,
+		    NULL);
+		printf("Device Name    : %s\n",str);
+
+		clGetDeviceInfo(devices[i], CL_DEVICE_VERSION, sizeof(str), str,
+		    NULL);
+		printf("Device Version : %s\n",str);
+
+		clGetDeviceInfo(devices[i], CL_DEVICE_GLOBAL_MEM_SIZE,
+		    sizeof(int_info), &int_info, NULL);
+		printf("Size of global memory : %lu (MB) \n",
+		    int_info/1024/1024);
+
+		clGetDeviceInfo(devices[i], CL_DEVICE_LOCAL_MEM_SIZE,
+		    sizeof(int_info), &int_info, NULL);
+		printf("Size of local memory  : %lu (KB) \n", int_info/1024);
+
+		clGetDeviceInfo(devices[i], CL_DEVICE_MAX_CLOCK_FREQUENCY,
+		    sizeof(int_info), &int_info, NULL);
+		printf("Max clock frequency   : %4.2lf (GHz) \n",
+		    int_info/1024.0);
+
+		printf("-------------------------------------------------------"
+		    "----\n\n");
+	}
+#endif
+
+	/*
+	 * XXX: Very AMD-centric, should make it more flexible...
+	 * Intel has the first device as the CPU and no GPU support
+	 */
+
+	/*
+	 * GPU
+	 */
+	puctx->gpu_queue = clCreateCommandQueueWithProperties(ctx, devices[0], 0,
+	    &result);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "GPU: clGetContextInfo() failed!");
+		goto err;
+	}
+
+	/*
+	 * CPU
+	 */
+#if 0
+	puctx->cpu_queue = clCreateCommandQueue(ctx, devices[1], 0,
+	    &result);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "CPU: clGetContextInfo() failed!");
+		goto err;
+	}
+#endif
+
+	puctx->devices = devices;
+	puctx->ctx = ctx;
+
+	free(platforms);
+
+	return CL_SUCCESS;
+
+err:
+	if (platforms)
+		free(platforms);
+	if (devices)
+		free(devices);
+	if (puctx->cpu_queue)
+		clReleaseCommandQueue(puctx->cpu_queue);
+	if (puctx->gpu_queue)
+		clReleaseCommandQueue(puctx->gpu_queue);
+	if (ctx)
+		clReleaseContext(ctx);
+
+
+	return result;
+}
+
+void
+cl_clean_up(struct cl_uctx uctx)
+{
+	if (uctx.devices) {
+		free(uctx.devices);
+	}
+	if (uctx.gpu_queue) {
+		clReleaseCommandQueue(uctx.gpu_queue);
+	}
+	if (uctx.cpu_queue) {
+		clReleaseCommandQueue(uctx.cpu_queue);
+	}
+	if (uctx.ctx) {
+		clReleaseContext(uctx.ctx);
+	}
+}
+
+cl_int
+cl_build(struct cl_uctx uctx, cl_device_type dev,
+    char *kern_fname, cl_program *pprogram)
+{
+	cl_program program = NULL;
+	cl_int result = CL_SUCCESS;
+
+	FILE *kern_file = NULL;
+	char *kern_src = NULL;
+	size_t srcsz = 0;
+
+	int type;
+
+#ifdef INTEL_KERNEL_DEBUG
+	char build_options[100] = "-g -s F:\\obj\\vs\\debug\\Bin\\Debug\\";
+	strcat(&build_options[32], kern_fname);
+#else
+	char *build_options = NULL;
+#endif
+
+	if (kern_fname == NULL || pprogram == NULL)
+		return CL_INVALID_VALUE;
+
+	/* XXX: AMD-centric, should probably be passed as param */
+	/* Decide the target based on device type */
+	if (dev == CL_DEVICE_TYPE_GPU) {
+		type = 0;
+	} else {
+		type = 1;
+	}
+
+	/*
+	 * Compile and link the OpenCL kernel.
+	 */
+
+	/* Read-in the source code */
+	kern_file = fopen(kern_fname, "rb");
+	if (kern_file == NULL) {
+		log_warn("cl", "Failed to open kernel source file %s!",
+		    kern_fname);
+		result = CL_INVALID_VALUE;
+		goto err;
+	}
+	fseek(kern_file, 0, SEEK_END);
+	srcsz = ftell(kern_file);
+	fseek(kern_file, 0, SEEK_SET);
+	kern_src = (char *)malloc(srcsz + 1);
+	if (kern_src == NULL) {
+		log_warn("cl", "kern_src malloc() failed!");
+		result = CL_INVALID_VALUE;
+		goto err;
+	}
+	fread(kern_src, 1, srcsz, kern_file);
+	kern_src[srcsz] = 0;
+
+	log_info("cl", "FILE DUMP BEGINS");
+	log_info("cl", "%s", kern_src);
+	log_info("cl", "FILE DUMP ENDS");
+
+	program = clCreateProgramWithSource(uctx.ctx, 1,
+	    (const char **)&kern_src, &srcsz, &result);
+	if (result != CL_SUCCESS) {
+		log_warn("cl", "clCreateProgamWithSource() failed!");
+		goto err;
+	}
+
+	/* Build the kernel */
+	result = clBuildProgram(program, 1, &uctx.devices[type],
+	    build_options, NULL, NULL);
+	if (result != CL_SUCCESS) {
+		/* Print out the build log in case of failure */
+		char programLog[10000] = {0};
+		log_warn("cl", "clBuildProgram() failed!");
+		clGetProgramBuildInfo(program, uctx.devices[type],
+		    CL_PROGRAM_BUILD_LOG, 10000, programLog, 0);
+		log_warn("cl", "%s\n", programLog);
+		goto err;
+	}
+
+	*pprogram = program;
+
+err:
+	if (kern_file)
+		fclose(kern_file);
+	if (kern_src)
+		free(kern_src);
+	if (result != CL_SUCCESS && program)
+		clReleaseProgram(program);
+	return result;
+}
+
+cl_int
+cl_get_kern(cl_program program, char *kname, cl_kernel *pkern)
+{
+	cl_int result = CL_SUCCESS;
+
+	if (program == NULL || kname == NULL || pkern == NULL)
+		return CL_INVALID_VALUE;
+
+	*pkern = clCreateKernel(program, kname, &result);
+
+#if 0
+	clGetKernelWorkGroupInfo(ckKernel[0], cdDevice[did],
+	    CL_KERNEL_WORK_GROUP_SIZE, sizeof(int_info), &int_info, NULL);
+	printf("GPU Maximum Work group size : %d\n", int_info);
+	clGetKernelWorkGroupInfo(ckKernel[0], cdDevice[did],
+	    CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(int_info),
+	    &int_info, NULL);
+	printf("GPU Preferred Work group size : %d\n", int_info);
+#endif
+
+	return result;
+}
diff --git a/cl_setup.h b/cl_setup.h
new file mode 100644
index 0000000..dce2b04
--- /dev/null
+++ b/cl_setup.h
@@ -0,0 +1,35 @@
+#ifndef SETUP_H__
+#define SETUP_H__
+
+#define CL_TARGET_OPENCL_VERSION 220	/* For Sinaia cluster */
+
+#include <CL/cl.h>
+
+/*
+ * OpenCL user context.
+ */
+struct cl_uctx {
+	cl_context ctx;				/* internal context */
+	cl_device_id *devices;			/* device in use */
+
+	cl_command_queue gpu_queue;		/* command queue */
+	cl_command_queue cpu_queue;		/* command queue */
+
+	cl_device_type device_type;		/* desired device type */
+	char platform_name[100];		/* desired platform */
+
+	size_t reduce_gws, reduce_lws;	/* norm global and local workspace */
+
+	/* BLAS kernels */
+	cl_kernel gemm_NN_kernel, gemm_TN_kernel;
+	cl_kernel frob_stage1_kernel, frob_stage2_kernel;
+};
+
+cl_int cl_init(struct cl_uctx *uctx);
+void cl_clean_up(struct cl_uctx uctx);
+
+cl_int cl_build(struct cl_uctx uctx, cl_device_type dev, char *kern_name,
+    cl_program *pprogram);
+cl_int cl_get_kern(cl_program program, char *kname, cl_kernel *pkern);
+
+#endif
diff --git a/clccubes.h b/clccubes.h
new file mode 100644
index 0000000..9234044
--- /dev/null
+++ b/clccubes.h
@@ -0,0 +1,26 @@
+#ifndef CLccubes_H__
+#define CLccubes_H__
+
+#include "cl_setup.h"
+
+#ifndef _MSC_VER
+#include <stdint.h>
+#else
+#include <stdint_msvc.h>
+#endif
+
+#define ROW_DIM 0
+#define COL_DIM 1
+
+struct ccubes_context {
+	struct cl_uctx *clctx;
+
+	cl_kernel ccubes_task;
+};
+
+int
+clccubes(struct ccubes_context *ccubesctx, cl_mem alpha0, cl_mem G, size_t signals,
+    size_t atoms, uint32_t sparsity, uint32_t pcoding, cl_mem gamma, cl_uint
+    num_events_in_wait_list, const cl_event *event_wait_list, cl_event *ev_ccubes);
+
+#endif
diff --git a/config.c b/config.c
new file mode 100644
index 0000000..42c299f
--- /dev/null
+++ b/config.c
@@ -0,0 +1,136 @@
+/*
+ * Copyright (c) 2012, Marius Barbu <msb@avengis.com>
+ * Copyright (c) 2013, Paul Irofti <paul@irofti.net>
+ */
+
+#include <stdio.h> /* sscanf/snprintf */
+#include <string.h> /* stricmp */
+#include <stdlib.h>
+
+#include "tree.h"
+#include "config.h"
+
+#if _MSC_VER
+#define snprintf _snprintf
+#endif
+
+struct ConfigMap {
+	char *name;
+	char *value;
+	RB_ENTRY(ConfigMap) entry;
+};
+
+int
+cfg_cmp(struct ConfigMap *cfg1, struct ConfigMap *cfg2)
+{
+	return (strcmp(cfg1->name, cfg2->name));
+}
+
+RB_HEAD(cfgtree, ConfigMap) cfg_head = RB_INITIALIZER(&cfg_head);
+RB_GENERATE(cfgtree, ConfigMap, entry, cfg_cmp);
+
+void
+config_set_string(char *name, char *value)
+{
+	struct ConfigMap *dup = NULL;
+	struct ConfigMap *cfg = malloc(sizeof *cfg);
+
+	if (cfg == NULL || name == NULL || value == NULL)
+		return;
+
+	cfg->name = strdup(name);
+	if (cfg->name == NULL)
+		return;
+	cfg->value = strdup(value);
+	if (cfg->value == NULL)
+		return;
+	if ((dup = RB_INSERT(cfgtree, &cfg_head, cfg))) {
+		dup->value = cfg->value;
+		free(cfg->name);
+		free(cfg);
+	}
+}
+
+void
+config_set_int(char *name, int value)
+{
+	char buff[20];
+
+	snprintf(buff, sizeof buff, "%d", value);
+
+	config_set_string(name, buff);
+}
+
+void
+config_set_bool(char *name, int value)
+{
+	char buff[20];
+
+	snprintf(buff, sizeof buff, "%d", value ? 1 : 0);
+
+	config_set_string(name, buff);
+}
+
+char*
+config_get_string(char *name, char *def)
+{
+	struct ConfigMap find, *found = NULL;
+	find.name = (char *)name;
+
+	found = RB_FIND(cfgtree, &cfg_head, &find);
+	if (found == NULL) {
+		return def;
+	}
+
+	return found->value;
+}
+
+int
+config_get_int(char *name, int def)
+{
+	int val = def;
+	struct ConfigMap find, *found = NULL;
+
+	find.name = (char *)name;
+	found = RB_FIND(cfgtree, &cfg_head, &find);
+	if (found == NULL) {
+		return def;
+	}
+
+	sscanf(found->value, "%i", &val);
+
+	return val;
+}
+
+int
+config_get_bool(char *name, int def)
+{
+	struct ConfigMap find, *found = NULL;
+
+	find.name = (char *)name;
+	found = RB_FIND(cfgtree, &cfg_head, &find);
+	if (found == NULL) {
+		return def;
+	}
+
+	if (strcmp(found->value, "0") == 0 ||
+	    strcmp(found->value, "false") == 0) {
+		return 0;
+	}
+
+	return 1;
+}
+
+void
+config_destroy()
+{
+	struct ConfigMap *var, *nxt;
+	for (var = RB_MIN(cfgtree, &cfg_head); var != NULL; var = nxt) {
+		nxt = RB_NEXT(cfgtree, &cfg_head, var);
+		RB_REMOVE(cfgtree, &cfg_head, var);
+		free(var->name);
+		free(var->value);
+		free(var);
+	}
+}
+
diff --git a/config.h b/config.h
new file mode 100644
index 0000000..89f169f
--- /dev/null
+++ b/config.h
@@ -0,0 +1,18 @@
+/*
+ * Copyright (c) 2012, Marius Barbu <msb@avengis.com>
+ * Copyright (c) 2013, Paul Irofti <paul@irofti.net>
+ */
+#ifndef CONFIG_H__
+#define CONFIG_H__
+
+void config_set_string(char *name, char *value);
+void config_set_int(char *name, int value);
+void config_set_bool(char *name, int value);
+
+char* config_get_string(char *name, char *def);
+int config_get_int(char *name, int def);
+int config_get_bool(char *name, int def);
+
+void config_destroy();
+
+#endif /* CONFIG_H__ */
diff --git a/logging.c b/logging.c
new file mode 100644
index 0000000..86b678c
--- /dev/null
+++ b/logging.c
@@ -0,0 +1,126 @@
+/*
+ * Copyright (c) 2012, Marius Barbu <msb@avengis.com>
+ * Copyright (c) 2013, Paul Irofti <paul@irofti.net>
+ */
+
+#include <stdio.h> /* FILE */
+#include <stdlib.h> /* abort */
+#include <stdarg.h>
+#include <string.h> /* strrchr */
+
+#include "tree.h"
+
+#include "config.h"
+#include "logging.h"
+
+static char *levelName[] = {
+	"INFO",
+	"DEBUG",
+	"WARN",
+	"ERROR"
+};
+
+static FILE *logFile;
+
+static char*
+nice_path(char *name)
+{
+	if (strrchr(name, '/')) {
+		name = strrchr(name, '/') + 1;
+	} else if (strrchr(name, '\\')) {
+		name = strrchr(name, '\\') + 1;
+	}
+
+	return name;
+}
+
+enum LogLevel
+log_level(char *id)
+{
+	char name[20] = { 0 };
+	size_t len = 0;
+	int cfgGlobal, cfgId;
+
+	if (id == NULL)
+		return LOG_LEVEL_ERROR;
+
+	len = strlen(id);
+	if (len == 0 || len > 15)
+		return LOG_LEVEL_ERROR;
+
+	strcpy(name, "log:");
+	strcat(name, id);
+	cfgGlobal = config_get_int("log", LOG_LEVEL_WARN);
+	cfgId = config_get_int(name, cfgGlobal);
+
+	if (cfgId > LOG_LEVEL_ERROR) {
+		cfgId = LOG_LEVEL_ERROR;
+	}
+
+	return cfgId;
+}
+
+static void
+log_message_raw_v(enum LogLevel l, char *id, char *fmt, va_list argp)
+{
+	if (l < log_level(id)) {
+		return;
+	}
+
+	if (logFile == 0) {
+		char *name = config_get_string("out", 0);
+		if (name) {
+			logFile = fopen(name, "a+t");
+			if (logFile == 0) {
+				perror(name);
+			}
+		}
+
+		if (logFile == 0) {
+			logFile = stdout;
+		}
+	}
+
+	vfprintf(logFile, fmt, argp);
+}
+
+void
+log_message(enum LogLevel l, char *where, int line, char *id,
+	char *fmt, ...)
+{
+	va_list argp;
+
+	if (l < log_level(id)) {
+		return;
+	}
+
+#ifdef NON_DIFFABLE_OUTPUT
+	log_message_raw(l, id, "%s - %s:%d - %s: ", levelName[l],
+	    nice_path(where), line, id);
+#endif
+	log_message_raw(l, id, "%s - %s: ", levelName[l], id);
+	va_start(argp, fmt);
+	log_message_raw_v(l, id, fmt, argp);
+	va_end(argp);
+	log_message_raw(l, id, "\n");
+	if (l >= LOG_LEVEL_ERROR) {
+		fflush(logFile);
+		/*exit(1);*/
+	}
+	fflush(logFile);
+}
+
+void
+log_message_raw(enum LogLevel l, char *id, char *fmt, ...)
+{	
+	va_list argp;
+	if (l < log_level(id)) {
+		return;
+	}
+
+
+	va_start(argp, fmt);
+	log_message_raw_v(l, id, fmt, argp);
+	va_end(argp);
+}
+
diff --git a/logging.h b/logging.h
new file mode 100644
index 0000000..99799f3
--- /dev/null
+++ b/logging.h
@@ -0,0 +1,74 @@
+/*
+ * Copyright (c) 2012, Marius Barbu <msb@avengis.com>
+ * Copyright (c) 2013, Paul Irofti <paul@irofti.net>
+ */
+#ifndef LOGGING_H__
+#define LOGGING_H__
+
+enum LogLevel {
+	LOG_LEVEL_INFO,
+	LOG_LEVEL_DEBUG,
+	LOG_LEVEL_WARN,
+	LOG_LEVEL_ERROR,
+	LOG_LEVEL_DISABLE,
+};
+
+enum LogLevel log_level(char *id);
+void log_message_raw(enum LogLevel l, char *id, char *fmt, ...);
+void log_message(enum LogLevel l, char *where, int line, char *id,
+	char *fmt, ...);
+
+#define log_info(id, ...) \
+	do { \
+	int enabled = log_level(id) <= LOG_LEVEL_INFO; \
+	if (enabled) { \
+		log_message(LOG_LEVEL_INFO, __FILE__, __LINE__, \
+			id, __VA_ARGS__); \
+	} \
+	} while (0)
+
+#define log_info_raw(id, ...) \
+	do { \
+	int enabled = log_level(id) <= LOG_LEVEL_INFO; \
+	if (enabled) { \
+		log_message_raw(LOG_LEVEL_INFO, \
+			id, __VA_ARGS__); \
+	} \
+	} while (0)
+
+#define log_debug(id, ...) \
+	do { \
+	int enabled = log_level(id) <= LOG_LEVEL_DEBUG; \
+	if (enabled) { \
+		log_message(LOG_LEVEL_DEBUG, __FILE__, __LINE__, \
+			id, __VA_ARGS__); \
+	} \
+	} while (0)
+
+#define log_debug_raw(id, ...) \
+	do { \
+	int enabled = log_level(id) <= LOG_LEVEL_DEBUG; \
+	if (enabled) { \
+		log_message_raw(LOG_LEVEL_DEBUG, \
+			id, __VA_ARGS__); \
+	} \
+	} while (0)
+
+
+#define log_warn(id, ...) \
+	log_message(LOG_LEVEL_WARN, __FILE__, __LINE__, \
+		id, __VA_ARGS__)
+		
+#define log_warn_raw(id, ...) \
+	log_message_raw(LOG_LEVEL_WARN, \
+		id, __VA_ARGS__)
+
+#define log_error(id, ...) \
+	log_message(LOG_LEVEL_ERROR, __FILE__, __LINE__, \
+		id, __VA_ARGS__)
+
+#define log_error_raw(id, ...) \
+	log_message(LOG_LEVEL_ERROR, \
+		id, __VA_ARGS__)
+
+#endif /* LOGGING_H__ */
diff --git a/queue.h b/queue.h
new file mode 100644
index 0000000..fb23a72
--- /dev/null
+++ b/queue.h
@@ -0,0 +1,568 @@
+/*	$OpenBSD: queue.h,v 1.35 2012/01/11 00:06:48 bluhm Exp $	*/
+/*	$NetBSD: queue.h,v 1.11 1996/05/16 05:17:14 mycroft Exp $	*/
+
+/*
+ * Copyright (c) 1991, 1993
+ *	The Regents of the University of California.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ * 1. Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ * 3. Neither the name of the University nor the names of its contributors
+ *    may be used to endorse or promote products derived from this software
+ *    without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE REGENTS AND CONTRIBUTORS ``AS IS'' AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED.  IN NO EVENT SHALL THE REGENTS OR CONTRIBUTORS BE LIABLE
+ * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+ * DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS
+ * OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION)
+ * HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
+ * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY
+ * OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF
+ * SUCH DAMAGE.
+ *
+ *	@(#)queue.h	8.5 (Berkeley) 8/20/94
+ */
+
+#ifndef	_SYS_QUEUE_H_
+#define	_SYS_QUEUE_H_
+
+/*
+ * This file defines five types of data structures: singly-linked lists, 
+ * lists, simple queues, tail queues, and circular queues.
+ *
+ *
+ * A singly-linked list is headed by a single forward pointer. The elements
+ * are singly linked for minimum space and pointer manipulation overhead at
+ * the expense of O(n) removal for arbitrary elements. New elements can be
+ * added to the list after an existing element or at the head of the list.
+ * Elements being removed from the head of the list should use the explicit
+ * macro for this purpose for optimum efficiency. A singly-linked list may
+ * only be traversed in the forward direction.  Singly-linked lists are ideal
+ * for applications with large datasets and few or no removals or for
+ * implementing a LIFO queue.
+ *
+ * A list is headed by a single forward pointer (or an array of forward
+ * pointers for a hash table header). The elements are doubly linked
+ * so that an arbitrary element can be removed without a need to
+ * traverse the list. New elements can be added to the list before
+ * or after an existing element or at the head of the list. A list
+ * may only be traversed in the forward direction.
+ *
+ * A simple queue is headed by a pair of pointers, one the head of the
+ * list and the other to the tail of the list. The elements are singly
+ * linked to save space, so elements can only be removed from the
+ * head of the list. New elements can be added to the list before or after
+ * an existing element, at the head of the list, or at the end of the
+ * list. A simple queue may only be traversed in the forward direction.
+ *
+ * A tail queue is headed by a pair of pointers, one to the head of the
+ * list and the other to the tail of the list. The elements are doubly
+ * linked so that an arbitrary element can be removed without a need to
+ * traverse the list. New elements can be added to the list before or
+ * after an existing element, at the head of the list, or at the end of
+ * the list. A tail queue may be traversed in either direction.
+ *
+ * A circle queue is headed by a pair of pointers, one to the head of the
+ * list and the other to the tail of the list. The elements are doubly
+ * linked so that an arbitrary element can be removed without a need to
+ * traverse the list. New elements can be added to the list before or after
+ * an existing element, at the head of the list, or at the end of the list.
+ * A circle queue may be traversed in either direction, but has a more
+ * complex end of list detection.
+ *
+ * For details on the use of these macros, see the queue(3) manual page.
+ */
+
+#if defined(QUEUE_MACRO_DEBUG) || (defined(_KERNEL) && defined(DIAGNOSTIC))
+#define _Q_INVALIDATE(a) (a) = ((void *)-1)
+#else
+#define _Q_INVALIDATE(a)
+#endif
+
+/*
+ * Singly-linked List definitions.
+ */
+#define SLIST_HEAD(name, type)						\
+struct name {								\
+	struct type *slh_first;	/* first element */			\
+}
+ 
+#define	SLIST_HEAD_INITIALIZER(head)					\
+	{ NULL }
+ 
+#define SLIST_ENTRY(type)						\
+struct {								\
+	struct type *sle_next;	/* next element */			\
+}
+ 
+/*
+ * Singly-linked List access methods.
+ */
+#define	SLIST_FIRST(head)	((head)->slh_first)
+#define	SLIST_END(head)		NULL
+#define	SLIST_EMPTY(head)	(SLIST_FIRST(head) == SLIST_END(head))
+#define	SLIST_NEXT(elm, field)	((elm)->field.sle_next)
+
+#define	SLIST_FOREACH(var, head, field)					\
+	for((var) = SLIST_FIRST(head);					\
+	    (var) != SLIST_END(head);					\
+	    (var) = SLIST_NEXT(var, field))
+
+#define	SLIST_FOREACH_SAFE(var, head, field, tvar)			\
+	for ((var) = SLIST_FIRST(head);				\
+	    (var) && ((tvar) = SLIST_NEXT(var, field), 1);		\
+	    (var) = (tvar))
+
+/*
+ * Singly-linked List functions.
+ */
+#define	SLIST_INIT(head) {						\
+	SLIST_FIRST(head) = SLIST_END(head);				\
+}
+
+#define	SLIST_INSERT_AFTER(slistelm, elm, field) do {			\
+	(elm)->field.sle_next = (slistelm)->field.sle_next;		\
+	(slistelm)->field.sle_next = (elm);				\
+} while (0)
+
+#define	SLIST_INSERT_HEAD(head, elm, field) do {			\
+	(elm)->field.sle_next = (head)->slh_first;			\
+	(head)->slh_first = (elm);					\
+} while (0)
+
+#define	SLIST_REMOVE_NEXT(head, elm, field) do {			\
+	(elm)->field.sle_next = (elm)->field.sle_next->field.sle_next;	\
+} while (0)
+
+#define	SLIST_REMOVE_HEAD(head, field) do {				\
+	(head)->slh_first = (head)->slh_first->field.sle_next;		\
+} while (0)
+
+#define SLIST_REMOVE(head, elm, type, field) do {			\
+	if ((head)->slh_first == (elm)) {				\
+		SLIST_REMOVE_HEAD((head), field);			\
+	} else {							\
+		struct type *curelm = (head)->slh_first;		\
+									\
+		while (curelm->field.sle_next != (elm))			\
+			curelm = curelm->field.sle_next;		\
+		curelm->field.sle_next =				\
+		    curelm->field.sle_next->field.sle_next;		\
+		_Q_INVALIDATE((elm)->field.sle_next);			\
+	}								\
+} while (0)
+
+/*
+ * List definitions.
+ */
+#define LIST_HEAD(name, type)						\
+struct name {								\
+	struct type *lh_first;	/* first element */			\
+}
+
+#define LIST_HEAD_INITIALIZER(head)					\
+	{ NULL }
+
+#define LIST_ENTRY(type)						\
+struct {								\
+	struct type *le_next;	/* next element */			\
+	struct type **le_prev;	/* address of previous next element */	\
+}
+
+/*
+ * List access methods
+ */
+#define	LIST_FIRST(head)		((head)->lh_first)
+#define	LIST_END(head)			NULL
+#define	LIST_EMPTY(head)		(LIST_FIRST(head) == LIST_END(head))
+#define	LIST_NEXT(elm, field)		((elm)->field.le_next)
+
+#define LIST_FOREACH(var, head, field)					\
+	for((var) = LIST_FIRST(head);					\
+	    (var)!= LIST_END(head);					\
+	    (var) = LIST_NEXT(var, field))
+
+#define	LIST_FOREACH_SAFE(var, head, field, tvar)			\
+	for ((var) = LIST_FIRST(head);				\
+	    (var) && ((tvar) = LIST_NEXT(var, field), 1);		\
+	    (var) = (tvar))
+
+/*
+ * List functions.
+ */
+#define	LIST_INIT(head) do {						\
+	LIST_FIRST(head) = LIST_END(head);				\
+} while (0)
+
+#define LIST_INSERT_AFTER(listelm, elm, field) do {			\
+	if (((elm)->field.le_next = (listelm)->field.le_next) != NULL)	\
+		(listelm)->field.le_next->field.le_prev =		\
+		    &(elm)->field.le_next;				\
+	(listelm)->field.le_next = (elm);				\
+	(elm)->field.le_prev = &(listelm)->field.le_next;		\
+} while (0)
+
+#define	LIST_INSERT_BEFORE(listelm, elm, field) do {			\
+	(elm)->field.le_prev = (listelm)->field.le_prev;		\
+	(elm)->field.le_next = (listelm);				\
+	*(listelm)->field.le_prev = (elm);				\
+	(listelm)->field.le_prev = &(elm)->field.le_next;		\
+} while (0)
+
+#define LIST_INSERT_HEAD(head, elm, field) do {				\
+	if (((elm)->field.le_next = (head)->lh_first) != NULL)		\
+		(head)->lh_first->field.le_prev = &(elm)->field.le_next;\
+	(head)->lh_first = (elm);					\
+	(elm)->field.le_prev = &(head)->lh_first;			\
+} while (0)
+
+#define LIST_REMOVE(elm, field) do {					\
+	if ((elm)->field.le_next != NULL)				\
+		(elm)->field.le_next->field.le_prev =			\
+		    (elm)->field.le_prev;				\
+	*(elm)->field.le_prev = (elm)->field.le_next;			\
+	_Q_INVALIDATE((elm)->field.le_prev);				\
+	_Q_INVALIDATE((elm)->field.le_next);				\
+} while (0)
+
+#define LIST_REPLACE(elm, elm2, field) do {				\
+	if (((elm2)->field.le_next = (elm)->field.le_next) != NULL)	\
+		(elm2)->field.le_next->field.le_prev =			\
+		    &(elm2)->field.le_next;				\
+	(elm2)->field.le_prev = (elm)->field.le_prev;			\
+	*(elm2)->field.le_prev = (elm2);				\
+	_Q_INVALIDATE((elm)->field.le_prev);				\
+	_Q_INVALIDATE((elm)->field.le_next);				\
+} while (0)
+
+/*
+ * Simple queue definitions.
+ */
+#define SIMPLEQ_HEAD(name, type)					\
+struct name {								\
+	struct type *sqh_first;	/* first element */			\
+	struct type **sqh_last;	/* addr of last next element */		\
+}
+
+#define SIMPLEQ_HEAD_INITIALIZER(head)					\
+	{ NULL, &(head).sqh_first }
+
+#define SIMPLEQ_ENTRY(type)						\
+struct {								\
+	struct type *sqe_next;	/* next element */			\
+}
+
+/*
+ * Simple queue access methods.
+ */
+#define	SIMPLEQ_FIRST(head)	    ((head)->sqh_first)
+#define	SIMPLEQ_END(head)	    NULL
+#define	SIMPLEQ_EMPTY(head)	    (SIMPLEQ_FIRST(head) == SIMPLEQ_END(head))
+#define	SIMPLEQ_NEXT(elm, field)    ((elm)->field.sqe_next)
+
+#define SIMPLEQ_FOREACH(var, head, field)				\
+	for((var) = SIMPLEQ_FIRST(head);				\
+	    (var) != SIMPLEQ_END(head);					\
+	    (var) = SIMPLEQ_NEXT(var, field))
+
+#define	SIMPLEQ_FOREACH_SAFE(var, head, field, tvar)			\
+	for ((var) = SIMPLEQ_FIRST(head);				\
+	    (var) && ((tvar) = SIMPLEQ_NEXT(var, field), 1);		\
+	    (var) = (tvar))
+
+/*
+ * Simple queue functions.
+ */
+#define	SIMPLEQ_INIT(head) do {						\
+	(head)->sqh_first = NULL;					\
+	(head)->sqh_last = &(head)->sqh_first;				\
+} while (0)
+
+#define SIMPLEQ_INSERT_HEAD(head, elm, field) do {			\
+	if (((elm)->field.sqe_next = (head)->sqh_first) == NULL)	\
+		(head)->sqh_last = &(elm)->field.sqe_next;		\
+	(head)->sqh_first = (elm);					\
+} while (0)
+
+#define SIMPLEQ_INSERT_TAIL(head, elm, field) do {			\
+	(elm)->field.sqe_next = NULL;					\
+	*(head)->sqh_last = (elm);					\
+	(head)->sqh_last = &(elm)->field.sqe_next;			\
+} while (0)
+
+#define SIMPLEQ_INSERT_AFTER(head, listelm, elm, field) do {		\
+	if (((elm)->field.sqe_next = (listelm)->field.sqe_next) == NULL)\
+		(head)->sqh_last = &(elm)->field.sqe_next;		\
+	(listelm)->field.sqe_next = (elm);				\
+} while (0)
+
+#define SIMPLEQ_REMOVE_HEAD(head, field) do {			\
+	if (((head)->sqh_first = (head)->sqh_first->field.sqe_next) == NULL) \
+		(head)->sqh_last = &(head)->sqh_first;			\
+} while (0)
+
+#define SIMPLEQ_REMOVE_NEXT(head, elm, field) do {			\
+	if (((elm)->field.sqe_next = (elm)->field.sqe_next->field.sqe_next) \
+	    == NULL)							\
+		(head)->sqh_last = &(elm)->field.sqe_next;		\
+} while (0)
+
+/*
+ * Tail queue definitions.
+ */
+#define TAILQ_HEAD(name, type)						\
+struct name {								\
+	struct type *tqh_first;	/* first element */			\
+	struct type **tqh_last;	/* addr of last next element */		\
+}
+
+#define TAILQ_HEAD_INITIALIZER(head)					\
+	{ NULL, &(head).tqh_first }
+
+#define TAILQ_ENTRY(type)						\
+struct {								\
+	struct type *tqe_next;	/* next element */			\
+	struct type **tqe_prev;	/* address of previous next element */	\
+}
+
+/* 
+ * tail queue access methods 
+ */
+#define	TAILQ_FIRST(head)		((head)->tqh_first)
+#define	TAILQ_END(head)			NULL
+#define	TAILQ_NEXT(elm, field)		((elm)->field.tqe_next)
+#define TAILQ_LAST(head, headname)					\
+	(*(((struct headname *)((head)->tqh_last))->tqh_last))
+/* XXX */
+#define TAILQ_PREV(elm, headname, field)				\
+	(*(((struct headname *)((elm)->field.tqe_prev))->tqh_last))
+#define	TAILQ_EMPTY(head)						\
+	(TAILQ_FIRST(head) == TAILQ_END(head))
+
+#define TAILQ_FOREACH(var, head, field)					\
+	for((var) = TAILQ_FIRST(head);					\
+	    (var) != TAILQ_END(head);					\
+	    (var) = TAILQ_NEXT(var, field))
+
+#define	TAILQ_FOREACH_SAFE(var, head, field, tvar)			\
+	for ((var) = TAILQ_FIRST(head);					\
+	    (var) != TAILQ_END(head) &&					\
+	    ((tvar) = TAILQ_NEXT(var, field), 1);			\
+	    (var) = (tvar))
+
+
+#define TAILQ_FOREACH_REVERSE(var, head, headname, field)		\
+	for((var) = TAILQ_LAST(head, headname);				\
+	    (var) != TAILQ_END(head);					\
+	    (var) = TAILQ_PREV(var, headname, field))
+
+#define	TAILQ_FOREACH_REVERSE_SAFE(var, head, headname, field, tvar)	\
+	for ((var) = TAILQ_LAST(head, headname);			\
+	    (var) != TAILQ_END(head) &&					\
+	    ((tvar) = TAILQ_PREV(var, headname, field), 1);		\
+	    (var) = (tvar))
+
+/*
+ * Tail queue functions.
+ */
+#define	TAILQ_INIT(head) do {						\
+	(head)->tqh_first = NULL;					\
+	(head)->tqh_last = &(head)->tqh_first;				\
+} while (0)
+
+#define TAILQ_INSERT_HEAD(head, elm, field) do {			\
+	if (((elm)->field.tqe_next = (head)->tqh_first) != NULL)	\
+		(head)->tqh_first->field.tqe_prev =			\
+		    &(elm)->field.tqe_next;				\
+	else								\
+		(head)->tqh_last = &(elm)->field.tqe_next;		\
+	(head)->tqh_first = (elm);					\
+	(elm)->field.tqe_prev = &(head)->tqh_first;			\
+} while (0)
+
+#define TAILQ_INSERT_TAIL(head, elm, field) do {			\
+	(elm)->field.tqe_next = NULL;					\
+	(elm)->field.tqe_prev = (head)->tqh_last;			\
+	*(head)->tqh_last = (elm);					\
+	(head)->tqh_last = &(elm)->field.tqe_next;			\
+} while (0)
+
+#define TAILQ_INSERT_AFTER(head, listelm, elm, field) do {		\
+	if (((elm)->field.tqe_next = (listelm)->field.tqe_next) != NULL)\
+		(elm)->field.tqe_next->field.tqe_prev =			\
+		    &(elm)->field.tqe_next;				\
+	else								\
+		(head)->tqh_last = &(elm)->field.tqe_next;		\
+	(listelm)->field.tqe_next = (elm);				\
+	(elm)->field.tqe_prev = &(listelm)->field.tqe_next;		\
+} while (0)
+
+#define	TAILQ_INSERT_BEFORE(listelm, elm, field) do {			\
+	(elm)->field.tqe_prev = (listelm)->field.tqe_prev;		\
+	(elm)->field.tqe_next = (listelm);				\
+	*(listelm)->field.tqe_prev = (elm);				\
+	(listelm)->field.tqe_prev = &(elm)->field.tqe_next;		\
+} while (0)
+
+#define TAILQ_REMOVE(head, elm, field) do {				\
+	if (((elm)->field.tqe_next) != NULL)				\
+		(elm)->field.tqe_next->field.tqe_prev =			\
+		    (elm)->field.tqe_prev;				\
+	else								\
+		(head)->tqh_last = (elm)->field.tqe_prev;		\
+	*(elm)->field.tqe_prev = (elm)->field.tqe_next;			\
+	_Q_INVALIDATE((elm)->field.tqe_prev);				\
+	_Q_INVALIDATE((elm)->field.tqe_next);				\
+} while (0)
+
+#define TAILQ_REPLACE(head, elm, elm2, field) do {			\
+	if (((elm2)->field.tqe_next = (elm)->field.tqe_next) != NULL)	\
+		(elm2)->field.tqe_next->field.tqe_prev =		\
+		    &(elm2)->field.tqe_next;				\
+	else								\
+		(head)->tqh_last = &(elm2)->field.tqe_next;		\
+	(elm2)->field.tqe_prev = (elm)->field.tqe_prev;			\
+	*(elm2)->field.tqe_prev = (elm2);				\
+	_Q_INVALIDATE((elm)->field.tqe_prev);				\
+	_Q_INVALIDATE((elm)->field.tqe_next);				\
+} while (0)
+
+/*
+ * Circular queue definitions.
+ */
+#define CIRCLEQ_HEAD(name, type)					\
+struct name {								\
+	struct type *cqh_first;		/* first element */		\
+	struct type *cqh_last;		/* last element */		\
+}
+
+#define CIRCLEQ_HEAD_INITIALIZER(head)					\
+	{ CIRCLEQ_END(&head), CIRCLEQ_END(&head) }
+
+#define CIRCLEQ_ENTRY(type)						\
+struct {								\
+	struct type *cqe_next;		/* next element */		\
+	struct type *cqe_prev;		/* previous element */		\
+}
+
+/*
+ * Circular queue access methods 
+ */
+#define	CIRCLEQ_FIRST(head)		((head)->cqh_first)
+#define	CIRCLEQ_LAST(head)		((head)->cqh_last)
+#define	CIRCLEQ_END(head)		((void *)(head))
+#define	CIRCLEQ_NEXT(elm, field)	((elm)->field.cqe_next)
+#define	CIRCLEQ_PREV(elm, field)	((elm)->field.cqe_prev)
+#define	CIRCLEQ_EMPTY(head)						\
+	(CIRCLEQ_FIRST(head) == CIRCLEQ_END(head))
+
+#define CIRCLEQ_FOREACH(var, head, field)				\
+	for((var) = CIRCLEQ_FIRST(head);				\
+	    (var) != CIRCLEQ_END(head);					\
+	    (var) = CIRCLEQ_NEXT(var, field))
+
+#define	CIRCLEQ_FOREACH_SAFE(var, head, field, tvar)			\
+	for ((var) = CIRCLEQ_FIRST(head);				\
+	    (var) != CIRCLEQ_END(head) &&				\
+	    ((tvar) = CIRCLEQ_NEXT(var, field), 1);			\
+	    (var) = (tvar))
+
+#define CIRCLEQ_FOREACH_REVERSE(var, head, field)			\
+	for((var) = CIRCLEQ_LAST(head);					\
+	    (var) != CIRCLEQ_END(head);					\
+	    (var) = CIRCLEQ_PREV(var, field))
+
+#define	CIRCLEQ_FOREACH_REVERSE_SAFE(var, head, headname, field, tvar)	\
+	for ((var) = CIRCLEQ_LAST(head, headname);			\
+	    (var) != CIRCLEQ_END(head) && 				\
+	    ((tvar) = CIRCLEQ_PREV(var, headname, field), 1);		\
+	    (var) = (tvar))
+
+/*
+ * Circular queue functions.
+ */
+#define	CIRCLEQ_INIT(head) do {						\
+	(head)->cqh_first = CIRCLEQ_END(head);				\
+	(head)->cqh_last = CIRCLEQ_END(head);				\
+} while (0)
+
+#define CIRCLEQ_INSERT_AFTER(head, listelm, elm, field) do {		\
+	(elm)->field.cqe_next = (listelm)->field.cqe_next;		\
+	(elm)->field.cqe_prev = (listelm);				\
+	if ((listelm)->field.cqe_next == CIRCLEQ_END(head))		\
+		(head)->cqh_last = (elm);				\
+	else								\
+		(listelm)->field.cqe_next->field.cqe_prev = (elm);	\
+	(listelm)->field.cqe_next = (elm);				\
+} while (0)
+
+#define CIRCLEQ_INSERT_BEFORE(head, listelm, elm, field) do {		\
+	(elm)->field.cqe_next = (listelm);				\
+	(elm)->field.cqe_prev = (listelm)->field.cqe_prev;		\
+	if ((listelm)->field.cqe_prev == CIRCLEQ_END(head))		\
+		(head)->cqh_first = (elm);				\
+	else								\
+		(listelm)->field.cqe_prev->field.cqe_next = (elm);	\
+	(listelm)->field.cqe_prev = (elm);				\
+} while (0)
+
+#define CIRCLEQ_INSERT_HEAD(head, elm, field) do {			\
+	(elm)->field.cqe_next = (head)->cqh_first;			\
+	(elm)->field.cqe_prev = CIRCLEQ_END(head);			\
+	if ((head)->cqh_last == CIRCLEQ_END(head))			\
+		(head)->cqh_last = (elm);				\
+	else								\
+		(head)->cqh_first->field.cqe_prev = (elm);		\
+	(head)->cqh_first = (elm);					\
+} while (0)
+
+#define CIRCLEQ_INSERT_TAIL(head, elm, field) do {			\
+	(elm)->field.cqe_next = CIRCLEQ_END(head);			\
+	(elm)->field.cqe_prev = (head)->cqh_last;			\
+	if ((head)->cqh_first == CIRCLEQ_END(head))			\
+		(head)->cqh_first = (elm);				\
+	else								\
+		(head)->cqh_last->field.cqe_next = (elm);		\
+	(head)->cqh_last = (elm);					\
+} while (0)
+
+#define	CIRCLEQ_REMOVE(head, elm, field) do {				\
+	if ((elm)->field.cqe_next == CIRCLEQ_END(head))			\
+		(head)->cqh_last = (elm)->field.cqe_prev;		\
+	else								\
+		(elm)->field.cqe_next->field.cqe_prev =			\
+		    (elm)->field.cqe_prev;				\
+	if ((elm)->field.cqe_prev == CIRCLEQ_END(head))			\
+		(head)->cqh_first = (elm)->field.cqe_next;		\
+	else								\
+		(elm)->field.cqe_prev->field.cqe_next =			\
+		    (elm)->field.cqe_next;				\
+	_Q_INVALIDATE((elm)->field.cqe_prev);				\
+	_Q_INVALIDATE((elm)->field.cqe_next);				\
+} while (0)
+
+#define CIRCLEQ_REPLACE(head, elm, elm2, field) do {			\
+	if (((elm2)->field.cqe_next = (elm)->field.cqe_next) ==		\
+	    CIRCLEQ_END(head))						\
+		(head).cqh_last = (elm2);				\
+	else								\
+		(elm2)->field.cqe_next->field.cqe_prev = (elm2);	\
+	if (((elm2)->field.cqe_prev = (elm)->field.cqe_prev) ==		\
+	    CIRCLEQ_END(head))						\
+		(head).cqh_first = (elm2);				\
+	else								\
+		(elm2)->field.cqe_prev->field.cqe_next = (elm2);	\
+	_Q_INVALIDATE((elm)->field.cqe_prev);				\
+	_Q_INVALIDATE((elm)->field.cqe_next);				\
+} while (0)
+
+#endif	/* !_SYS_QUEUE_H_ */
diff --git a/real.h b/real.h
new file mode 100644
index 0000000..56069c7
--- /dev/null
+++ b/real.h
@@ -0,0 +1,20 @@
+#ifndef REAL_H__
+#define REAL_H__
+
+#include <math.h>
+
+#ifdef USE_DOUBLE
+typedef double real;
+#define R_ZERO	1e-8
+#define rabs fabs
+#define rsqrt sqrt
+#define clAmdBlasGemmEx clAmdBlasDgemmEx
+#else
+typedef float real;
+#define R_ZERO	1e-6
+#define rabs fabsf
+#define rsqrt sqrtf
+#define clAmdBlasGemmEx clAmdBlasSgemmEx
+#endif
+
+#endif
diff --git a/test_ccubes.c b/test_ccubes.c
new file mode 100644
index 0000000..66f8d22
--- /dev/null
+++ b/test_ccubes.c
@@ -0,0 +1,140 @@
+/*
+ * Copyright (c) 2013, 2014, Paul Irofti <paul@irofti.net>
+ */
+
+#include <stdio.h>
+#include <string.h>
+
+#include "real.h"
+#include "cl_setup.h"
+
+#include "clccubes.h"
+
+#include "config.h"
+#include "logging.h"
+
+int
+parse_args(char *argv[], int argc,
+    uint32_t *signals, uint32_t *atoms, uint32_t *problem_size,
+    uint32_t *sparsity, uint32_t *iters, uint32_t *patoms, uint32_t *pcoding,
+    real *rmse)
+{
+	int i = 1;
+	while (i < argc) {
+		if (strcmp(argv[i], "-signals") == 0) {
+			*signals = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-atoms") == 0) {
+			*atoms = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-problem") == 0) {
+			*problem_size = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-sparsity") == 0) {
+			*sparsity = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-iterations") == 0) {
+			*iters = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-patoms") == 0) {
+			*patoms = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-pcoding") == 0) {
+			*pcoding = strtoul(argv[++i], NULL, 0);
+		} else if (strcmp(argv[i], "-rmse") == 0) {
+			sscanf(argv[++i], "%f", rmse);
+		} else {
+			printf("Unrecognized option %s\n", argv[i]);
+			return 1;
+		}
+		i++;
+	}
+
+	return 0;
+}
+
+void
+usage(char *program)
+{
+	printf("USAGE: %s [-signals sigs] [-atoms atoms] "
+	    "[-problem problem_size] [-sparsity s] [-iterations iters] "
+	    "[-patoms p] [-pcoding c]\n",
+	    program);
+	exit(1);
+}
+
+int main(int argc, char *argv[])
+{
+	int rc = 0;
+
+	real *zeroes = NULL;
+	uint64_t start, stop;
+	char log[100] = {0};
+	char *program;
+
+	struct ccubes_context ccubesctx;
+	cl_program ccubes_program;
+
+
+	/* Inputs */
+
+	/* Outputs */
+
+	/* Rounding */
+
+	/* Logging */
+	config_set_int("log", LOG_LEVEL_WARN);
+	config_set_int("log:ccubes", LOG_LEVEL_WARN);
+	config_set_int("log:test", LOG_LEVEL_WARN);
+	config_set_int("cl", LOG_LEVEL_DEBUG);
+
+#if 0
+	rc = parse_args(argv, argc,
+	    &signals, &atoms, &problem_size, &sparsity, &iters,
+	    &patoms, &pcoding, &target_error);
+	if (rc != 0) {
+		usage(argv[0]);
+	}
+
+	if (strstr(argv[0], "test_ccubes") != NULL)
+		program = "ccubes";
+	sprintf(log, "log-%s-p%d-n%d-m%d-k%d-i%d-pa%d", program,
+	    problem_size, atoms, signals, sparsity, iters, patoms);
+	config_set_string("out", log);
+#endif
+
+	log_debug("test", "Starting ccubes...");
+
+	/*
+	 *	CCubes
+	 */
+	ccubesctx.clctx = malloc(sizeof *ccubesctx.clctx);
+
+	ccubesctx.clctx->device_type = CL_DEVICE_TYPE_GPU;
+	strcpy(ccubesctx.clctx->platform_name, "NVIDIA Corporation\0");
+	/* strcpy(ccubesctx.clctx->platform_name, "Advanced Micro Devices, Inc.\0"); */
+	/*strcpy(ccubesctx.clctx->platform_name, "Intel(R) Corporation\0");*/
+	rc = cl_init(ccubesctx.clctx);
+	if (rc != CL_SUCCESS) {
+		printf("[%d] Failed to initialize the OpenCL framework\n",
+		    rc);
+		goto err;
+	}
+
+
+	rc = cl_build(*ccubesctx.clctx, CL_DEVICE_TYPE_GPU,
+	    "ccubes.cl", &ccubes_program);
+	if (rc != CL_SUCCESS) {
+		log_warn("test", "Failed building ccubes.cl (%d)", rc);
+		goto err;
+	}
+
+	rc = cl_get_kern(ccubes_program, "ccubes_task",
+	    &ccubesctx.ccubes_task);
+	if (rc != CL_SUCCESS) {
+		log_warn("test", "Failed fetching ccubes_task (%d)", rc);
+		goto err;
+	}
+
+	fflush(stdout);
+err:
+	clReleaseProgram(ccubes_program);
+
+	cl_clean_up(*ccubesctx.clctx);
+
+	return rc;
+}
diff --git a/tree.h b/tree.h
new file mode 100644
index 0000000..80d0f53
--- /dev/null
+++ b/tree.h
@@ -0,0 +1,748 @@
+/*	$OpenBSD: tree.h,v 1.13 2011/07/09 00:19:45 pirofti Exp $	*/
+/*
+ * Copyright 2002 Niels Provos <provos@citi.umich.edu>
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ * 1. Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE AUTHOR ``AS IS'' AND ANY EXPRESS OR
+ * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES
+ * OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
+ * IN NO EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT,
+ * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT
+ * NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+ * DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+ * THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
+ * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef	_SYS_TREE_H_
+#define	_SYS_TREE_H_
+
+/*
+ * This file defines data structures for different types of trees:
+ * splay trees and red-black trees.
+ *
+ * A splay tree is a self-organizing data structure.  Every operation
+ * on the tree causes a splay to happen.  The splay moves the requested
+ * node to the root of the tree and partly rebalances it.
+ *
+ * This has the benefit that request locality causes faster lookups as
+ * the requested nodes move to the top of the tree.  On the other hand,
+ * every lookup causes memory writes.
+ *
+ * The Balance Theorem bounds the total access time for m operations
+ * and n inserts on an initially empty tree as O((m + n)lg n).  The
+ * amortized cost for a sequence of m accesses to a splay tree is O(lg n);
+ *
+ * A red-black tree is a binary search tree with the node color as an
+ * extra attribute.  It fulfills a set of conditions:
+ *	- every search path from the root to a leaf consists of the
+ *	  same number of black nodes,
+ *	- each red node (except for the root) has a black parent,
+ *	- each leaf node is black.
+ *
+ * Every operation on a red-black tree is bounded as O(lg n).
+ * The maximum height of a red-black tree is 2lg (n+1).
+ */
+
+#define SPLAY_HEAD(name, type)						\
+struct name {								\
+	struct type *sph_root; /* root of the tree */			\
+}
+
+#define SPLAY_INITIALIZER(root)						\
+	{ NULL }
+
+#define SPLAY_INIT(root) do {						\
+	(root)->sph_root = NULL;					\
+} while (0)
+
+#define SPLAY_ENTRY(type)						\
+struct {								\
+	struct type *spe_left; /* left element */			\
+	struct type *spe_right; /* right element */			\
+}
+
+#define SPLAY_LEFT(elm, field)		(elm)->field.spe_left
+#define SPLAY_RIGHT(elm, field)		(elm)->field.spe_right
+#define SPLAY_ROOT(head)		(head)->sph_root
+#define SPLAY_EMPTY(head)		(SPLAY_ROOT(head) == NULL)
+
+/* SPLAY_ROTATE_{LEFT,RIGHT} expect that tmp hold SPLAY_{RIGHT,LEFT} */
+#define SPLAY_ROTATE_RIGHT(head, tmp, field) do {			\
+	SPLAY_LEFT((head)->sph_root, field) = SPLAY_RIGHT(tmp, field);	\
+	SPLAY_RIGHT(tmp, field) = (head)->sph_root;			\
+	(head)->sph_root = tmp;						\
+} while (0)
+	
+#define SPLAY_ROTATE_LEFT(head, tmp, field) do {			\
+	SPLAY_RIGHT((head)->sph_root, field) = SPLAY_LEFT(tmp, field);	\
+	SPLAY_LEFT(tmp, field) = (head)->sph_root;			\
+	(head)->sph_root = tmp;						\
+} while (0)
+
+#define SPLAY_LINKLEFT(head, tmp, field) do {				\
+	SPLAY_LEFT(tmp, field) = (head)->sph_root;			\
+	tmp = (head)->sph_root;						\
+	(head)->sph_root = SPLAY_LEFT((head)->sph_root, field);		\
+} while (0)
+
+#define SPLAY_LINKRIGHT(head, tmp, field) do {				\
+	SPLAY_RIGHT(tmp, field) = (head)->sph_root;			\
+	tmp = (head)->sph_root;						\
+	(head)->sph_root = SPLAY_RIGHT((head)->sph_root, field);	\
+} while (0)
+
+#define SPLAY_ASSEMBLE(head, node, left, right, field) do {		\
+	SPLAY_RIGHT(left, field) = SPLAY_LEFT((head)->sph_root, field);	\
+	SPLAY_LEFT(right, field) = SPLAY_RIGHT((head)->sph_root, field);\
+	SPLAY_LEFT((head)->sph_root, field) = SPLAY_RIGHT(node, field);	\
+	SPLAY_RIGHT((head)->sph_root, field) = SPLAY_LEFT(node, field);	\
+} while (0)
+
+/* Generates prototypes and inline functions */
+
+#define SPLAY_PROTOTYPE(name, type, field, cmp)				\
+void name##_SPLAY(struct name *, struct type *);			\
+void name##_SPLAY_MINMAX(struct name *, int);				\
+struct type *name##_SPLAY_INSERT(struct name *, struct type *);		\
+struct type *name##_SPLAY_REMOVE(struct name *, struct type *);		\
+									\
+/* Finds the node with the same key as elm */				\
+static __inline struct type *						\
+name##_SPLAY_FIND(struct name *head, struct type *elm)			\
+{									\
+	if (SPLAY_EMPTY(head))						\
+		return(NULL);						\
+	name##_SPLAY(head, elm);					\
+	if ((cmp)(elm, (head)->sph_root) == 0)				\
+		return (head->sph_root);				\
+	return (NULL);							\
+}									\
+									\
+static __inline struct type *						\
+name##_SPLAY_NEXT(struct name *head, struct type *elm)			\
+{									\
+	name##_SPLAY(head, elm);					\
+	if (SPLAY_RIGHT(elm, field) != NULL) {				\
+		elm = SPLAY_RIGHT(elm, field);				\
+		while (SPLAY_LEFT(elm, field) != NULL) {		\
+			elm = SPLAY_LEFT(elm, field);			\
+		}							\
+	} else								\
+		elm = NULL;						\
+	return (elm);							\
+}									\
+									\
+static __inline struct type *						\
+name##_SPLAY_MIN_MAX(struct name *head, int val)			\
+{									\
+	name##_SPLAY_MINMAX(head, val);					\
+        return (SPLAY_ROOT(head));					\
+}
+
+/* Main splay operation.
+ * Moves node close to the key of elm to top
+ */
+#define SPLAY_GENERATE(name, type, field, cmp)				\
+struct type *								\
+name##_SPLAY_INSERT(struct name *head, struct type *elm)		\
+{									\
+    if (SPLAY_EMPTY(head)) {						\
+	    SPLAY_LEFT(elm, field) = SPLAY_RIGHT(elm, field) = NULL;	\
+    } else {								\
+	    int __comp;							\
+	    name##_SPLAY(head, elm);					\
+	    __comp = (cmp)(elm, (head)->sph_root);			\
+	    if(__comp < 0) {						\
+		    SPLAY_LEFT(elm, field) = SPLAY_LEFT((head)->sph_root, field);\
+		    SPLAY_RIGHT(elm, field) = (head)->sph_root;		\
+		    SPLAY_LEFT((head)->sph_root, field) = NULL;		\
+	    } else if (__comp > 0) {					\
+		    SPLAY_RIGHT(elm, field) = SPLAY_RIGHT((head)->sph_root, field);\
+		    SPLAY_LEFT(elm, field) = (head)->sph_root;		\
+		    SPLAY_RIGHT((head)->sph_root, field) = NULL;	\
+	    } else							\
+		    return ((head)->sph_root);				\
+    }									\
+    (head)->sph_root = (elm);						\
+    return (NULL);							\
+}									\
+									\
+struct type *								\
+name##_SPLAY_REMOVE(struct name *head, struct type *elm)		\
+{									\
+	struct type *__tmp;						\
+	if (SPLAY_EMPTY(head))						\
+		return (NULL);						\
+	name##_SPLAY(head, elm);					\
+	if ((cmp)(elm, (head)->sph_root) == 0) {			\
+		if (SPLAY_LEFT((head)->sph_root, field) == NULL) {	\
+			(head)->sph_root = SPLAY_RIGHT((head)->sph_root, field);\
+		} else {						\
+			__tmp = SPLAY_RIGHT((head)->sph_root, field);	\
+			(head)->sph_root = SPLAY_LEFT((head)->sph_root, field);\
+			name##_SPLAY(head, elm);			\
+			SPLAY_RIGHT((head)->sph_root, field) = __tmp;	\
+		}							\
+		return (elm);						\
+	}								\
+	return (NULL);							\
+}									\
+									\
+void									\
+name##_SPLAY(struct name *head, struct type *elm)			\
+{									\
+	struct type __node, *__left, *__right, *__tmp;			\
+	int __comp;							\
+\
+	SPLAY_LEFT(&__node, field) = SPLAY_RIGHT(&__node, field) = NULL;\
+	__left = __right = &__node;					\
+\
+	while ((__comp = (cmp)(elm, (head)->sph_root))) {		\
+		if (__comp < 0) {					\
+			__tmp = SPLAY_LEFT((head)->sph_root, field);	\
+			if (__tmp == NULL)				\
+				break;					\
+			if ((cmp)(elm, __tmp) < 0){			\
+				SPLAY_ROTATE_RIGHT(head, __tmp, field);	\
+				if (SPLAY_LEFT((head)->sph_root, field) == NULL)\
+					break;				\
+			}						\
+			SPLAY_LINKLEFT(head, __right, field);		\
+		} else if (__comp > 0) {				\
+			__tmp = SPLAY_RIGHT((head)->sph_root, field);	\
+			if (__tmp == NULL)				\
+				break;					\
+			if ((cmp)(elm, __tmp) > 0){			\
+				SPLAY_ROTATE_LEFT(head, __tmp, field);	\
+				if (SPLAY_RIGHT((head)->sph_root, field) == NULL)\
+					break;				\
+			}						\
+			SPLAY_LINKRIGHT(head, __left, field);		\
+		}							\
+	}								\
+	SPLAY_ASSEMBLE(head, &__node, __left, __right, field);		\
+}									\
+									\
+/* Splay with either the minimum or the maximum element			\
+ * Used to find minimum or maximum element in tree.			\
+ */									\
+void name##_SPLAY_MINMAX(struct name *head, int __comp) \
+{									\
+	struct type __node, *__left, *__right, *__tmp;			\
+\
+	SPLAY_LEFT(&__node, field) = SPLAY_RIGHT(&__node, field) = NULL;\
+	__left = __right = &__node;					\
+\
+	while (1) {							\
+		if (__comp < 0) {					\
+			__tmp = SPLAY_LEFT((head)->sph_root, field);	\
+			if (__tmp == NULL)				\
+				break;					\
+			if (__comp < 0){				\
+				SPLAY_ROTATE_RIGHT(head, __tmp, field);	\
+				if (SPLAY_LEFT((head)->sph_root, field) == NULL)\
+					break;				\
+			}						\
+			SPLAY_LINKLEFT(head, __right, field);		\
+		} else if (__comp > 0) {				\
+			__tmp = SPLAY_RIGHT((head)->sph_root, field);	\
+			if (__tmp == NULL)				\
+				break;					\
+			if (__comp > 0) {				\
+				SPLAY_ROTATE_LEFT(head, __tmp, field);	\
+				if (SPLAY_RIGHT((head)->sph_root, field) == NULL)\
+					break;				\
+			}						\
+			SPLAY_LINKRIGHT(head, __left, field);		\
+		}							\
+	}								\
+	SPLAY_ASSEMBLE(head, &__node, __left, __right, field);		\
+}
+
+#define SPLAY_NEGINF	-1
+#define SPLAY_INF	1
+
+#define SPLAY_INSERT(name, x, y)	name##_SPLAY_INSERT(x, y)
+#define SPLAY_REMOVE(name, x, y)	name##_SPLAY_REMOVE(x, y)
+#define SPLAY_FIND(name, x, y)		name##_SPLAY_FIND(x, y)
+#define SPLAY_NEXT(name, x, y)		name##_SPLAY_NEXT(x, y)
+#define SPLAY_MIN(name, x)		(SPLAY_EMPTY(x) ? NULL	\
+					: name##_SPLAY_MIN_MAX(x, SPLAY_NEGINF))
+#define SPLAY_MAX(name, x)		(SPLAY_EMPTY(x) ? NULL	\
+					: name##_SPLAY_MIN_MAX(x, SPLAY_INF))
+
+#define SPLAY_FOREACH(x, name, head)					\
+	for ((x) = SPLAY_MIN(name, head);				\
+	     (x) != NULL;						\
+	     (x) = SPLAY_NEXT(name, head, x))
+
+/* Macros that define a red-black tree */
+#define RB_HEAD(name, type)						\
+struct name {								\
+	struct type *rbh_root; /* root of the tree */			\
+}
+
+#define RB_INITIALIZER(root)						\
+	{ NULL }
+
+#define RB_INIT(root) do {						\
+	(root)->rbh_root = NULL;					\
+} while (0)
+
+#define RB_BLACK	0
+#define RB_RED		1
+#define RB_ENTRY(type)							\
+struct {								\
+	struct type *rbe_left;		/* left element */		\
+	struct type *rbe_right;		/* right element */		\
+	struct type *rbe_parent;	/* parent element */		\
+	int rbe_color;			/* node color */		\
+}
+
+#define RB_LEFT(elm, field)		(elm)->field.rbe_left
+#define RB_RIGHT(elm, field)		(elm)->field.rbe_right
+#define RB_PARENT(elm, field)		(elm)->field.rbe_parent
+#define RB_COLOR(elm, field)		(elm)->field.rbe_color
+#define RB_ROOT(head)			(head)->rbh_root
+#define RB_EMPTY(head)			(RB_ROOT(head) == NULL)
+
+#define RB_SET(elm, parent, field) do {					\
+	RB_PARENT(elm, field) = parent;					\
+	RB_LEFT(elm, field) = RB_RIGHT(elm, field) = NULL;		\
+	RB_COLOR(elm, field) = RB_RED;					\
+} while (0)
+
+#define RB_SET_BLACKRED(black, red, field) do {				\
+	RB_COLOR(black, field) = RB_BLACK;				\
+	RB_COLOR(red, field) = RB_RED;					\
+} while (0)
+
+#ifndef RB_AUGMENT
+#define RB_AUGMENT(x)	do {} while (0)
+#endif
+
+#define RB_ROTATE_LEFT(head, elm, tmp, field) do {			\
+	(tmp) = RB_RIGHT(elm, field);					\
+	if ((RB_RIGHT(elm, field) = RB_LEFT(tmp, field))) {		\
+		RB_PARENT(RB_LEFT(tmp, field), field) = (elm);		\
+	}								\
+	RB_AUGMENT(elm);						\
+	if ((RB_PARENT(tmp, field) = RB_PARENT(elm, field))) {		\
+		if ((elm) == RB_LEFT(RB_PARENT(elm, field), field))	\
+			RB_LEFT(RB_PARENT(elm, field), field) = (tmp);	\
+		else							\
+			RB_RIGHT(RB_PARENT(elm, field), field) = (tmp);	\
+	} else								\
+		(head)->rbh_root = (tmp);				\
+	RB_LEFT(tmp, field) = (elm);					\
+	RB_PARENT(elm, field) = (tmp);					\
+	RB_AUGMENT(tmp);						\
+	if ((RB_PARENT(tmp, field)))					\
+		RB_AUGMENT(RB_PARENT(tmp, field));			\
+} while (0)
+
+#define RB_ROTATE_RIGHT(head, elm, tmp, field) do {			\
+	(tmp) = RB_LEFT(elm, field);					\
+	if ((RB_LEFT(elm, field) = RB_RIGHT(tmp, field))) {		\
+		RB_PARENT(RB_RIGHT(tmp, field), field) = (elm);		\
+	}								\
+	RB_AUGMENT(elm);						\
+	if ((RB_PARENT(tmp, field) = RB_PARENT(elm, field))) {		\
+		if ((elm) == RB_LEFT(RB_PARENT(elm, field), field))	\
+			RB_LEFT(RB_PARENT(elm, field), field) = (tmp);	\
+		else							\
+			RB_RIGHT(RB_PARENT(elm, field), field) = (tmp);	\
+	} else								\
+		(head)->rbh_root = (tmp);				\
+	RB_RIGHT(tmp, field) = (elm);					\
+	RB_PARENT(elm, field) = (tmp);					\
+	RB_AUGMENT(tmp);						\
+	if ((RB_PARENT(tmp, field)))					\
+		RB_AUGMENT(RB_PARENT(tmp, field));			\
+} while (0)
+
+/* Generates prototypes and inline functions */
+#define	RB_PROTOTYPE(name, type, field, cmp)				\
+	RB_PROTOTYPE_INTERNAL(name, type, field, cmp,)
+#define	RB_PROTOTYPE_STATIC(name, type, field, cmp)			\
+	RB_PROTOTYPE_INTERNAL(name, type, field, cmp, __attribute__((__unused__)) static)
+#define RB_PROTOTYPE_INTERNAL(name, type, field, cmp, attr)		\
+attr void name##_RB_INSERT_COLOR(struct name *, struct type *);		\
+attr void name##_RB_REMOVE_COLOR(struct name *, struct type *, struct type *);\
+attr struct type *name##_RB_REMOVE(struct name *, struct type *);	\
+attr struct type *name##_RB_INSERT(struct name *, struct type *);	\
+attr struct type *name##_RB_FIND(struct name *, struct type *);		\
+attr struct type *name##_RB_NFIND(struct name *, struct type *);	\
+attr struct type *name##_RB_NEXT(struct type *);			\
+attr struct type *name##_RB_PREV(struct type *);			\
+attr struct type *name##_RB_MINMAX(struct name *, int);			\
+									\
+
+/* Main rb operation.
+ * Moves node close to the key of elm to top
+ */
+#define	RB_GENERATE(name, type, field, cmp)				\
+	RB_GENERATE_INTERNAL(name, type, field, cmp,)
+#define	RB_GENERATE_STATIC(name, type, field, cmp)			\
+	RB_GENERATE_INTERNAL(name, type, field, cmp, __attribute__((__unused__)) static)
+#define RB_GENERATE_INTERNAL(name, type, field, cmp, attr)		\
+attr void								\
+name##_RB_INSERT_COLOR(struct name *head, struct type *elm)		\
+{									\
+	struct type *parent, *gparent, *tmp;				\
+	while ((parent = RB_PARENT(elm, field)) &&			\
+	    RB_COLOR(parent, field) == RB_RED) {			\
+		gparent = RB_PARENT(parent, field);			\
+		if (parent == RB_LEFT(gparent, field)) {		\
+			tmp = RB_RIGHT(gparent, field);			\
+			if (tmp && RB_COLOR(tmp, field) == RB_RED) {	\
+				RB_COLOR(tmp, field) = RB_BLACK;	\
+				RB_SET_BLACKRED(parent, gparent, field);\
+				elm = gparent;				\
+				continue;				\
+			}						\
+			if (RB_RIGHT(parent, field) == elm) {		\
+				RB_ROTATE_LEFT(head, parent, tmp, field);\
+				tmp = parent;				\
+				parent = elm;				\
+				elm = tmp;				\
+			}						\
+			RB_SET_BLACKRED(parent, gparent, field);	\
+			RB_ROTATE_RIGHT(head, gparent, tmp, field);	\
+		} else {						\
+			tmp = RB_LEFT(gparent, field);			\
+			if (tmp && RB_COLOR(tmp, field) == RB_RED) {	\
+				RB_COLOR(tmp, field) = RB_BLACK;	\
+				RB_SET_BLACKRED(parent, gparent, field);\
+				elm = gparent;				\
+				continue;				\
+			}						\
+			if (RB_LEFT(parent, field) == elm) {		\
+				RB_ROTATE_RIGHT(head, parent, tmp, field);\
+				tmp = parent;				\
+				parent = elm;				\
+				elm = tmp;				\
+			}						\
+			RB_SET_BLACKRED(parent, gparent, field);	\
+			RB_ROTATE_LEFT(head, gparent, tmp, field);	\
+		}							\
+	}								\
+	RB_COLOR(head->rbh_root, field) = RB_BLACK;			\
+}									\
+									\
+attr void								\
+name##_RB_REMOVE_COLOR(struct name *head, struct type *parent, struct type *elm) \
+{									\
+	struct type *tmp;						\
+	while ((elm == NULL || RB_COLOR(elm, field) == RB_BLACK) &&	\
+	    elm != RB_ROOT(head)) {					\
+		if (RB_LEFT(parent, field) == elm) {			\
+			tmp = RB_RIGHT(parent, field);			\
+			if (RB_COLOR(tmp, field) == RB_RED) {		\
+				RB_SET_BLACKRED(tmp, parent, field);	\
+				RB_ROTATE_LEFT(head, parent, tmp, field);\
+				tmp = RB_RIGHT(parent, field);		\
+			}						\
+			if ((RB_LEFT(tmp, field) == NULL ||		\
+			    RB_COLOR(RB_LEFT(tmp, field), field) == RB_BLACK) &&\
+			    (RB_RIGHT(tmp, field) == NULL ||		\
+			    RB_COLOR(RB_RIGHT(tmp, field), field) == RB_BLACK)) {\
+				RB_COLOR(tmp, field) = RB_RED;		\
+				elm = parent;				\
+				parent = RB_PARENT(elm, field);		\
+			} else {					\
+				if (RB_RIGHT(tmp, field) == NULL ||	\
+				    RB_COLOR(RB_RIGHT(tmp, field), field) == RB_BLACK) {\
+					struct type *oleft;		\
+					if ((oleft = RB_LEFT(tmp, field)))\
+						RB_COLOR(oleft, field) = RB_BLACK;\
+					RB_COLOR(tmp, field) = RB_RED;	\
+					RB_ROTATE_RIGHT(head, tmp, oleft, field);\
+					tmp = RB_RIGHT(parent, field);	\
+				}					\
+				RB_COLOR(tmp, field) = RB_COLOR(parent, field);\
+				RB_COLOR(parent, field) = RB_BLACK;	\
+				if (RB_RIGHT(tmp, field))		\
+					RB_COLOR(RB_RIGHT(tmp, field), field) = RB_BLACK;\
+				RB_ROTATE_LEFT(head, parent, tmp, field);\
+				elm = RB_ROOT(head);			\
+				break;					\
+			}						\
+		} else {						\
+			tmp = RB_LEFT(parent, field);			\
+			if (RB_COLOR(tmp, field) == RB_RED) {		\
+				RB_SET_BLACKRED(tmp, parent, field);	\
+				RB_ROTATE_RIGHT(head, parent, tmp, field);\
+				tmp = RB_LEFT(parent, field);		\
+			}						\
+			if ((RB_LEFT(tmp, field) == NULL ||		\
+			    RB_COLOR(RB_LEFT(tmp, field), field) == RB_BLACK) &&\
+			    (RB_RIGHT(tmp, field) == NULL ||		\
+			    RB_COLOR(RB_RIGHT(tmp, field), field) == RB_BLACK)) {\
+				RB_COLOR(tmp, field) = RB_RED;		\
+				elm = parent;				\
+				parent = RB_PARENT(elm, field);		\
+			} else {					\
+				if (RB_LEFT(tmp, field) == NULL ||	\
+				    RB_COLOR(RB_LEFT(tmp, field), field) == RB_BLACK) {\
+					struct type *oright;		\
+					if ((oright = RB_RIGHT(tmp, field)))\
+						RB_COLOR(oright, field) = RB_BLACK;\
+					RB_COLOR(tmp, field) = RB_RED;	\
+					RB_ROTATE_LEFT(head, tmp, oright, field);\
+					tmp = RB_LEFT(parent, field);	\
+				}					\
+				RB_COLOR(tmp, field) = RB_COLOR(parent, field);\
+				RB_COLOR(parent, field) = RB_BLACK;	\
+				if (RB_LEFT(tmp, field))		\
+					RB_COLOR(RB_LEFT(tmp, field), field) = RB_BLACK;\
+				RB_ROTATE_RIGHT(head, parent, tmp, field);\
+				elm = RB_ROOT(head);			\
+				break;					\
+			}						\
+		}							\
+	}								\
+	if (elm)							\
+		RB_COLOR(elm, field) = RB_BLACK;			\
+}									\
+									\
+attr struct type *							\
+name##_RB_REMOVE(struct name *head, struct type *elm)			\
+{									\
+	struct type *child, *parent, *old = elm;			\
+	int color;							\
+	if (RB_LEFT(elm, field) == NULL)				\
+		child = RB_RIGHT(elm, field);				\
+	else if (RB_RIGHT(elm, field) == NULL)				\
+		child = RB_LEFT(elm, field);				\
+	else {								\
+		struct type *left;					\
+		elm = RB_RIGHT(elm, field);				\
+		while ((left = RB_LEFT(elm, field)))			\
+			elm = left;					\
+		child = RB_RIGHT(elm, field);				\
+		parent = RB_PARENT(elm, field);				\
+		color = RB_COLOR(elm, field);				\
+		if (child)						\
+			RB_PARENT(child, field) = parent;		\
+		if (parent) {						\
+			if (RB_LEFT(parent, field) == elm)		\
+				RB_LEFT(parent, field) = child;		\
+			else						\
+				RB_RIGHT(parent, field) = child;	\
+			RB_AUGMENT(parent);				\
+		} else							\
+			RB_ROOT(head) = child;				\
+		if (RB_PARENT(elm, field) == old)			\
+			parent = elm;					\
+		(elm)->field = (old)->field;				\
+		if (RB_PARENT(old, field)) {				\
+			if (RB_LEFT(RB_PARENT(old, field), field) == old)\
+				RB_LEFT(RB_PARENT(old, field), field) = elm;\
+			else						\
+				RB_RIGHT(RB_PARENT(old, field), field) = elm;\
+			RB_AUGMENT(RB_PARENT(old, field));		\
+		} else							\
+			RB_ROOT(head) = elm;				\
+		RB_PARENT(RB_LEFT(old, field), field) = elm;		\
+		if (RB_RIGHT(old, field))				\
+			RB_PARENT(RB_RIGHT(old, field), field) = elm;	\
+		if (parent) {						\
+			left = parent;					\
+			do {						\
+				RB_AUGMENT(left);			\
+			} while ((left = RB_PARENT(left, field)));	\
+		}							\
+		goto color;						\
+	}								\
+	parent = RB_PARENT(elm, field);					\
+	color = RB_COLOR(elm, field);					\
+	if (child)							\
+		RB_PARENT(child, field) = parent;			\
+	if (parent) {							\
+		if (RB_LEFT(parent, field) == elm)			\
+			RB_LEFT(parent, field) = child;			\
+		else							\
+			RB_RIGHT(parent, field) = child;		\
+		RB_AUGMENT(parent);					\
+	} else								\
+		RB_ROOT(head) = child;					\
+color:									\
+	if (color == RB_BLACK)						\
+		name##_RB_REMOVE_COLOR(head, parent, child);		\
+	return (old);							\
+}									\
+									\
+/* Inserts a node into the RB tree */					\
+attr struct type *							\
+name##_RB_INSERT(struct name *head, struct type *elm)			\
+{									\
+	struct type *tmp;						\
+	struct type *parent = NULL;					\
+	int comp = 0;							\
+	tmp = RB_ROOT(head);						\
+	while (tmp) {							\
+		parent = tmp;						\
+		comp = (cmp)(elm, parent);				\
+		if (comp < 0)						\
+			tmp = RB_LEFT(tmp, field);			\
+		else if (comp > 0)					\
+			tmp = RB_RIGHT(tmp, field);			\
+		else							\
+			return (tmp);					\
+	}								\
+	RB_SET(elm, parent, field);					\
+	if (parent != NULL) {						\
+		if (comp < 0)						\
+			RB_LEFT(parent, field) = elm;			\
+		else							\
+			RB_RIGHT(parent, field) = elm;			\
+		RB_AUGMENT(parent);					\
+	} else								\
+		RB_ROOT(head) = elm;					\
+	name##_RB_INSERT_COLOR(head, elm);				\
+	return (NULL);							\
+}									\
+									\
+/* Finds the node with the same key as elm */				\
+attr struct type *							\
+name##_RB_FIND(struct name *head, struct type *elm)			\
+{									\
+	struct type *tmp = RB_ROOT(head);				\
+	int comp;							\
+	while (tmp) {							\
+		comp = cmp(elm, tmp);					\
+		if (comp < 0)						\
+			tmp = RB_LEFT(tmp, field);			\
+		else if (comp > 0)					\
+			tmp = RB_RIGHT(tmp, field);			\
+		else							\
+			return (tmp);					\
+	}								\
+	return (NULL);							\
+}									\
+									\
+/* Finds the first node greater than or equal to the search key */	\
+attr struct type *							\
+name##_RB_NFIND(struct name *head, struct type *elm)			\
+{									\
+	struct type *tmp = RB_ROOT(head);				\
+	struct type *res = NULL;					\
+	int comp;							\
+	while (tmp) {							\
+		comp = cmp(elm, tmp);					\
+		if (comp < 0) {						\
+			res = tmp;					\
+			tmp = RB_LEFT(tmp, field);			\
+		}							\
+		else if (comp > 0)					\
+			tmp = RB_RIGHT(tmp, field);			\
+		else							\
+			return (tmp);					\
+	}								\
+	return (res);							\
+}									\
+									\
+/* ARGSUSED */								\
+attr struct type *							\
+name##_RB_NEXT(struct type *elm)					\
+{									\
+	if (RB_RIGHT(elm, field)) {					\
+		elm = RB_RIGHT(elm, field);				\
+		while (RB_LEFT(elm, field))				\
+			elm = RB_LEFT(elm, field);			\
+	} else {							\
+		if (RB_PARENT(elm, field) &&				\
+		    (elm == RB_LEFT(RB_PARENT(elm, field), field)))	\
+			elm = RB_PARENT(elm, field);			\
+		else {							\
+			while (RB_PARENT(elm, field) &&			\
+			    (elm == RB_RIGHT(RB_PARENT(elm, field), field)))\
+				elm = RB_PARENT(elm, field);		\
+			elm = RB_PARENT(elm, field);			\
+		}							\
+	}								\
+	return (elm);							\
+}									\
+									\
+/* ARGSUSED */								\
+attr struct type *							\
+name##_RB_PREV(struct type *elm)					\
+{									\
+	if (RB_LEFT(elm, field)) {					\
+		elm = RB_LEFT(elm, field);				\
+		while (RB_RIGHT(elm, field))				\
+			elm = RB_RIGHT(elm, field);			\
+	} else {							\
+		if (RB_PARENT(elm, field) &&				\
+		    (elm == RB_RIGHT(RB_PARENT(elm, field), field)))	\
+			elm = RB_PARENT(elm, field);			\
+		else {							\
+			while (RB_PARENT(elm, field) &&			\
+			    (elm == RB_LEFT(RB_PARENT(elm, field), field)))\
+				elm = RB_PARENT(elm, field);		\
+			elm = RB_PARENT(elm, field);			\
+		}							\
+	}								\
+	return (elm);							\
+}									\
+									\
+attr struct type *							\
+name##_RB_MINMAX(struct name *head, int val)				\
+{									\
+	struct type *tmp = RB_ROOT(head);				\
+	struct type *parent = NULL;					\
+	while (tmp) {							\
+		parent = tmp;						\
+		if (val < 0)						\
+			tmp = RB_LEFT(tmp, field);			\
+		else							\
+			tmp = RB_RIGHT(tmp, field);			\
+	}								\
+	return (parent);						\
+}
+
+#define RB_NEGINF	-1
+#define RB_INF	1
+
+#define RB_INSERT(name, x, y)	name##_RB_INSERT(x, y)
+#define RB_REMOVE(name, x, y)	name##_RB_REMOVE(x, y)
+#define RB_FIND(name, x, y)	name##_RB_FIND(x, y)
+#define RB_NFIND(name, x, y)	name##_RB_NFIND(x, y)
+#define RB_NEXT(name, x, y)	name##_RB_NEXT(y)
+#define RB_PREV(name, x, y)	name##_RB_PREV(y)
+#define RB_MIN(name, x)		name##_RB_MINMAX(x, RB_NEGINF)
+#define RB_MAX(name, x)		name##_RB_MINMAX(x, RB_INF)
+
+#define RB_FOREACH(x, name, head)					\
+	for ((x) = RB_MIN(name, head);					\
+	     (x) != NULL;						\
+	     (x) = name##_RB_NEXT(x))
+
+#define RB_FOREACH_SAFE(x, name, head, y)				\
+	for ((x) = RB_MIN(name, head);					\
+	    ((x) != NULL) && ((y) = name##_RB_NEXT(x), 1);		\
+	     (x) = (y))
+
+#define RB_FOREACH_REVERSE(x, name, head)				\
+	for ((x) = RB_MAX(name, head);					\
+	     (x) != NULL;						\
+	     (x) = name##_RB_PREV(x))
+
+#define RB_FOREACH_REVERSE_SAFE(x, name, head, y)			\
+	for ((x) = RB_MAX(name, head);					\
+	    ((x) != NULL) && ((y) = name##_RB_PREV(x), 1);		\
+	     (x) = (y))
+
+#endif	/* _SYS_TREE_H_ */