 
 	// 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)));
+    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

    int blocks  = (numRows + threads - 1) / threads;

	T *			A;
	CUDA_TRY(cudaMalloc(&A, (sizeof(T) * csrReturn->nnz)));
	CUDA_TRY(cudaMemset(A, 0, (sizeof(T) * csrReturn->nnz)));

    // Now start moving the data and creating the CSR

    		status = runConverter<double>(gdfData, csrReturn, offsets);
    	    break;
    	default:
    		RMM_TRY(rmmFree(IA, 0));
    		RMM_TRY(rmmFree(JA, 0));
    		RMM_TRY(rmmFree(offsets, 0));
    		return GDF_UNSUPPORTED_DTYPE;
    }

    RMM_TRY(rmmFree(offsets, 0));

	return status;
}
