 }
 
 // copy the contents of a section of a ring buffer into a plain array
-__global__ 
-void ring_buffer_to_plain_array_kernel(
-    const uint32* buffer,
-    const uint32  buffer_size,
-    const uint32  begin,
-    const uint32  end,
-          uint32* output)
-{
-    const uint32 thread_id = threadIdx.x + BLOCKDIM*blockIdx.x;
-
-    if (begin + thread_id < end)
-        output[ thread_id ] = buffer[ (begin + thread_id) % buffer_size ];
-}
+struct ring_buffer_to_plain_array_functor
+{
+    // constructor
+    ring_buffer_to_plain_array_functor(
+        const uint32* _buffer,
+        const uint32  _buffer_size,
+        const uint32  _begin,
+        const uint32  _end,
+              uint32* _output)
+    : buffer( _buffer ),
+      buffer_size( _buffer_size ),
+      begin( _begin ),
+      end( _end ),
+      output( _output ) {}
+
+    // functor operator
+    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
+    void operator() (const uint32 i) const
+    {
+        if (begin + i < end)
+            output[i] = buffer[ (begin + i) % buffer_size ];
+    }
+
+    const uint32* buffer;
+    const uint32  buffer_size;
+    const uint32  begin;
+    const uint32  end;
+          uint32* output;
+};
 
 void ring_buffer_to_plain_array(
     const uint32* buffer,

 */

#include <nvBowtie/bowtie2/cuda/aligner.h>

namespace nvbio {
namespace bowtie2 {

    const uint32  end,
          uint32* output)
{
    nvbio::for_each<device_tag>(
        end - begin,
        thrust::make_counting_iterator<uint32>(0),
        ring_buffer_to_plain_array_functor(
            buffer,
            buffer_size,
            begin,
            end,
            output ) );
}

} // namespace cuda
