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

  extern __shared__ float buffer[];

  cudaDMAIndirect<true,false,ALIGNMENT>
    dma0(offsets,bytes_per_elmt,num_elmts);

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

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