 #include <gdf/gdf.h>
 #include <gdf/errorutils.h>
 #include "gdf/gdf_io.h"
+#include "rmm.h"
 
 #include <thrust/scan.h>
 #include <thrust/execution_policy.h>
+#include "thrust_rmm_allocator.h"
 
 using namespace std;
 

		determineValidRecCount<<<blocks, threads>>>(gdfData[x]->valid, numRows, numCol, offsets);
	}

	//--------------------------------------------------------------------------------------
	// Now do an exclusive scan to compute the offsets for where to write data
    thrust::exclusive_scan(thrust::device, offsets, (offsets + numRows + 1), offsets);

	//--------------------------------------------------------------------------------------
    // get the number of elements - NNZ, this is the last item in the array


	// Allocate space for the offset - this will eventually be IA - dtype is long since the sum of all column elements could be larger than int32
	gdf_size_type * offsets;
    RMM_TRY(rmmAlloc((void**)&offsets, (numRows + 2) * sizeof(int64_t), 0)); // TODO: non-default stream?
    CUDA_TRY(cudaMemset(offsets, 0, ( sizeof(int64_t) * (numRows + 2) ) ));

    // do a pass over each columns, and have each column updates the row count
