 static unsigned int
 compact
 (
-    const T             *in,
-    T                   *out,
-    const uint32_t      *flags,
+    const T             *d_in,
+    T                   *d_out,
+    const uint32_t      *d_flags,
     const uint32_t      length
 )
 {

    uint32_t blocks;

    permute_control(length, blocks, threads);
    permute_core< T,backward,compact >
        <<<blocks,threads>>>(in, out, indices, length, valid, num_valid);
}



    cudaMalloc((void**) &num, sizeof(uint32_t));
    cudaMalloc((void**) &indices, length * sizeof(uint32_t));

    if (backward) prescanr_plusui(d_flags, indices, length);
    else          prescanl_plusui(d_flags, indices, length);

    /*
     * At this point, we know exactly how many elements will be required for the
     * `out' array. Maybe we should allocate that array here?
     */
    permute<T,backward,true>(d_in, d_out, indices, length, d_flags, num);

    cudaMemcpy(&N, num, sizeof(uint32_t), cudaMemcpyDeviceToHost);
    cudaFree(num);
