 		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

#include <gdf/gdf.h>
#include <gdf/errorutils.h>
#include "gdf/gdf_io.h"

#include <thrust/scan.h>
#include <thrust/execution_policy.h>

using namespace std;


	//--------------------------------------------------------------------------------------
	// now start creating output data
    size_t * IA;
    RMM_TRY(rmmAlloc((void**)&IA, (numRows + 2) * sizeof(gdf_size_type), 0));
    CUDA_TRY(cudaMemcpy(IA, offsets, ( sizeof(gdf_size_type) * (numRows + 2) ), cudaMemcpyDeviceToDevice) );

    int64_t * 	JA;
    RMM_TRY( rmmAlloc((void**)&JA, (sizeof(int64_t) * nnz), 0));

    //----------------------------------------------------------------------------------
    // Now just missing A and the moving of data
