     const group_t *d_scratch_g,
     const unsigned int *d_scratch_idx,
     const unsigned int *d_offset,
-    group_t *d_pidx_group_table
+    group_t *d_pidx_group_table,
+    unsigned int pidx_group_table_pitch
     )
     {
     unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;


        // allocate temporary array
        unsigned int *d_offsets = (unsigned int *)alloc.allocate(group_size*n_groups*sizeof(unsigned int));
        thrust::device_ptr<unsigned int> offsets_ptr(d_offsets);

        // do a segmented prefix scan to determine dest offsets
        thrust::device_ptr<unsigned int> scratch_idx_ptr(d_scratch_idx);
        thrust::constant_iterator<unsigned int> const_it(pidx_group_table_pitch);
        thrust::exclusive_scan_by_key(thrust::cuda::par(alloc),
            scratch_idx_ptr,
            scratch_idx_ptr + group_size*n_groups,
            const_it,
            offsets_ptr);

        // scatter groups to destinations
        block_size = 512;

            d_scratch_g,
            d_scratch_idx,
            d_offsets,
            d_pidx_group_table,
            pidx_group_table_pitch);
 
        alloc.deallocate((char *) d_offsets,0);
        alloc.deallocate((char *) d_seg_offsets, 0);
        } 

    // release temporary arrays
