--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/bitonic_sort/cuda/swap.cu	Wed Mar 25 19:13:56 2015 +0900
@@ -0,0 +1,20 @@
+extern "C" {
+    __global__ void
+    swap(long* param, int* Data)
+    {
+        long block = param[0];
+        long first = param[1];
+        int x = blockIdx.x*blockDim.x+threadIdx.x;
+        
+        int position = x/block;
+        int index = x+block*position;
+        
+        block = (first == 1) ? ((block<<1)*(position+1))-(index%block)-1 : index+block;
+    
+        if (Data[block] < Data[index]) {
+            int tmp = Data[index];
+            Data[index] = Data[block];
+            Data[block] = tmp;
+        }
+    }        
+}

--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/example/bitonic_sort/cuda/task_init.cc	Wed Mar 25 19:13:56 2015 +0900
@@ -0,0 +1,9 @@
+#include "Func.h"
+#include "Scheduler.h"
+#include "CudaScheduler.h"
+
+void
+gpu_task_init(void)
+{
+    CudaSchedRegister(SWAP, "cuda/swap.ptx", "swap");
+}

--- a/example/bitonic_sort/main.cc	Mon Mar 23 18:31:58 2015 +0900
+++ b/example/bitonic_sort/main.cc	Wed Mar 25 19:13:56 2015 +0900
@@ -25,7 +25,7 @@
 int length = 1024;
 int* data;
 
-bool flag = false;
+bool wait_flag = false;
 
 static double
 getTime()
@@ -101,24 +101,27 @@
 
     HTask* wait;
     HTask* swap;
+    bool lastOne = false;
 
     for (int i=2; i <= length; i=2*i) {
-        bool first = true;
+        long first = 1;
+        lastOne = (length <= i*2) ? true : false;
         for (int j=i>>1; 0 < j; j=j>>1) {
             swap = manager->create_task(SWAP);
             swap->set_inData(0, data, length*sizeof(int*));
             swap->set_outData(0, data, length*sizeof(int*));
             swap->set_param(0, (long)j);
-            swap->set_param(1, (bool)first);
+            swap->set_param(1, (long)first);
             swap->set_cpu(spe_cpu);
-            swap->flip();
-            if (flag)
+            if (!(lastOne && !(0 < (j>>1))))
+                swap->flip();
+            if (wait_flag)
                 swap->wait_for(wait);
             swap->iterate(length/2);
-            first = false;
+            first = 0;
             wait = swap;
         }
-        flag = true;
+        wait_flag = true;
     }
 }
     
@@ -140,6 +143,6 @@
 {
     ed_time = getTime();
     printf("%0.6f\n",ed_time-st_time);
-    //print();
+    //    print();
     check_data();
 }

--- a/example/bitonic_sort/ppe/swap.cc	Mon Mar 23 18:31:58 2015 +0900
+++ b/example/bitonic_sort/ppe/swap.cc	Wed Mar 25 19:13:56 2015 +0900
@@ -8,13 +8,13 @@
     int* inData = (int*)s->get_input(rbuf, 0);
     int* outData = (int*)s->get_output(wbuf, 0);
     long block = (long)s->get_param(0);
-    bool first = (bool)s->get_param(1);
+    long first = (long)s->get_param(1);
     int x = s->x;
     
     int position = x/block;
     int index = x+block*position;
 
-    block = first ? ((block<<1)*(position+1))-(index%block)-1 : index+block;
+    block = (first == 1)? ((block<<1)*(position+1))-(index%block)-1 : index+block;
     
     if (inData[block] < inData[index]) {
         int tmp = inData[index];