             num_valid[0] = valid[length-1] + indices[length-1];
     }
 
-    if (idx < length)
+    if (lengthIsPow2 || idx < length)
     {
 	if (compact && valid[idx])
-	    out[indices[idx]] = in[idx];
+	    d_out[indices[idx]] = d_in[idx];
 	else if (backward)
-	    out[idx] = in[indices[idx]];
+	    d_out[idx] = d_in[indices[idx]];
 	else
-	    out[indices[idx]] = in[idx];
+	    d_out[indices[idx]] = d_in[idx];
     }
 }
 

 * 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>
__global__ static void
permute_core
(
    const T             *in,
    T                   *out,
    const uint32_t      *indices,
    const uint32_t      length,
    const uint32_t      *valid     = NULL,

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