     dma0 (1, num_dma_threads,
             num_compute_threads,
             num_compute_threads,
-            offsets,
             bytes_per_elmt,
             num_elmts);
 

    dma0.wait_for_dma_start();
    dma0.finish_async_dma();
#else
    dma0.execute_dma(&(buffer[ALIGN_OFFSET]), base_ptr);
#endif
  }
  else

  extern __shared__ float buffer[];

  cudaDMAIndirect<true,false,ALIGNMENT,BYTES_PER_ELMT,DMA_THREADS,NUM_ELMTS>
    dma0;

  zero_buffer(buffer, buffer_size);
  __syncthreads();
  float *base_ptr = &(idata[ALIGN_OFFSET]);
  dma0.execute_dma(offsets,base_ptr, &(buffer[ALIGN_OFFSET]));
  __syncthreads();
  copy_buffer(buffer, odata, buffer_size);
}
