 static void
 permute
 (
-    const T             *in,
-    T                   *out,
+    const T             *d_in,
+    T                   *d_out,
     const uint32_t      *indices,
     const uint32_t      length,
     const uint32_t      *valid     = NULL,

    uint32_t blocks;

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



 * We return the number of valid elements found via `num_valid', but blunder on
 * ahead regardless of whether the `out' array is large enough or not.
 */
template <typename T, bool backward, bool compact, bool lengthIsPow2>
__global__ static void
permute_core
(
    const T             *d_in,
    T                   *d_out,
    const uint32_t      *indices,
    const uint32_t      length,
    const uint32_t      *valid     = NULL,
