for (int32_t k = 0; k < 32; k++) {
    rsort_scan(cv.commands, cv.context,
        kernel_map[scan_name_str], kernel_map[update_name_str],
        g_in, g_zeros, 0, k, n);
    rsort_scan(cv.commands, cv.context,
        kernel_map[scan_name_str], kernel_map[update_name_str],
        g_in, g_ones, 1, k, n);

    err = clSetKernelArg(kernel_map[reassemble_name_str], 0, sizeof(cl_mem), &g_in);
    CHK_ERR(err);
    err = clSetKernelArg(kernel_map[reassemble_name_str], 1, sizeof(cl_mem), &g_out);
    CHK_ERR(err);
    err = clSetKernelArg(kernel_map[reassemble_name_str], 2, sizeof(cl_mem), &g_zeros);
    CHK_ERR(err);
    err = clSetKernelArg(kernel_map[reassemble_name_str], 3, sizeof(cl_mem), &g_ones);
    CHK_ERR(err);
    err = clSetKernelArg(kernel_map[reassemble_name_str], 4, sizeof(int), &k);
    CHK_ERR(err);
    err = clSetKernelArg(kernel_map[reassemble_name_str], 5, sizeof(int), &n);
    CHK_ERR(err);

    err = clEnqueueNDRangeKernel(cv.commands,
                     kernel_map[reassemble_name_str],
                     1,//work_dim,
                     NULL, //global_work_offset
                     global_work_size, //global_work_size
                     local_work_size, //local_work_size
                     0, //num_events_in_wait_list
                     NULL, //event_wait_list
                     NULL //
                     );
    CHK_ERR(err);

    cl_mem g_tPtr = g_in;
    g_in = g_out;
    g_out = g_tPtr;
  }