 		determineValidRecCount<<<blocks, threads>>>(gdfData[x]->valid, numRows, numCol, offsets);
 	}
 
+	rmm_temp_allocator allocator(0); // TODO: non-default stream?
+
 	//--------------------------------------------------------------------------------------
 	// Now do an exclusive scan to compute the offsets for where to write data
-    thrust::exclusive_scan(thrust::device, offsets, (offsets + numRows + 1), offsets);
+    thrust::exclusive_scan(thrust::cuda::par(allocator).on(0), 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;
    CUDA_TRY(cudaMalloc(&offsets, (numRows + 2) * sizeof(int64_t)));
    CUDA_TRY(cudaMemset(offsets, 0, ( sizeof(int64_t) * (numRows + 2) ) ));

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

#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;

